From 5a2575faca105cecbf44bc70f7f8a177f6c4f674 Mon Sep 17 00:00:00 2001 From: dirk Date: Fri, 15 Apr 2016 09:24:26 +0200 Subject: [PATCH] Complete rewrite of the OpenCL API. --- Magick++/lib/Functions.cpp | 20 +- Magick++/lib/Magick++/Functions.h | 2 +- MagickCore/accelerate-private.h | 120 +- MagickCore/accelerate.c | 1795 +++++------ MagickCore/accelerate.h | 3 +- MagickCore/magick.c | 2 + MagickCore/opencl-private.h | 640 ++-- MagickCore/opencl.c | 4870 +++++++++++++---------------- MagickCore/opencl.h | 63 +- 9 files changed, 3231 insertions(+), 4284 deletions(-) diff --git a/Magick++/lib/Functions.cpp b/Magick++/lib/Functions.cpp index 0db54c4f83..a4da42ef94 100644 --- a/Magick++/lib/Functions.cpp +++ b/Magick++/lib/Functions.cpp @@ -28,28 +28,16 @@ MagickPPExport void Magick::CloneString(char **destination_, MagickPPExport void Magick::DisableOpenCL(void) { - GetPPException; - MagickCore::InitImageMagickOpenCL(MagickCore::MAGICK_OPENCL_OFF,NULL,NULL, - exceptionInfo); - ThrowPPException(false); + MagickCore::SetOpenCLEnabled(MagickFalse); } -MagickPPExport bool Magick::EnableOpenCL(const bool useCache_) +MagickPPExport bool Magick::EnableOpenCL(void) { bool status; - GetPPException; - if (useCache_) - status=MagickCore::InitImageMagickOpenCL( - MagickCore::MAGICK_OPENCL_DEVICE_SELECT_AUTO,NULL,NULL,exceptionInfo) == - MagickTrue; - else - status=MagickCore::InitImageMagickOpenCL( - MagickCore::MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE,NULL,NULL, - exceptionInfo) == MagickTrue; - ThrowPPException(false); - return(status); + status=MagickCore::SetOpenCLEnabled(MagickTrue) != MagickFalse; + return(status); } MagickPPExport void Magick::InitializeMagick(const char *path_) diff --git a/Magick++/lib/Magick++/Functions.h b/Magick++/lib/Magick++/Functions.h index f109f9dc33..863cab1386 100644 --- a/Magick++/lib/Magick++/Functions.h +++ b/Magick++/lib/Magick++/Functions.h @@ -23,7 +23,7 @@ namespace Magick MagickPPExport void DisableOpenCL(void); // Enable OpenCL acceleration (only works when build with OpenCL support) - MagickPPExport bool EnableOpenCL(const bool useCache_=true); + MagickPPExport bool EnableOpenCL(void); // C library initialization routine MagickPPExport void InitializeMagick(const char *path_); diff --git a/MagickCore/accelerate-private.h b/MagickCore/accelerate-private.h index c72d08ee18..0b7db5d2e2 100644 --- a/MagickCore/accelerate-private.h +++ b/MagickCore/accelerate-private.h @@ -49,7 +49,7 @@ typedef struct _FloatPixelPacket black; } FloatPixelPacket; -const char* accelerateKernels = +const char *accelerateKernels = /* Define declarations. @@ -2387,79 +2387,6 @@ OPENCL_ENDIF() } ) -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -% R a n d o m % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -*/ - -STRINGIFY( - - inline float GetPseudoRandomValue(uint4* seed, const float normalizeRand) { - uint4 s = *seed; - do { - unsigned int alpha = (unsigned int)(s.y ^ (s.y << 11)); - s.y = s.z; - s.z = s.w; - s.w = s.x; - s.x = (s.x ^ (s.x >> 19)) ^ (alpha ^ (alpha >> 8)); - } while (s.x == ~0UL); - *seed = s; - return (normalizeRand*s.x); - } - - __kernel void RandomNumberGenerator(__global uint* seeds, const float normalizeRand - , __global float* randomNumbers, const uint init - , const uint numRandomNumbers) { - - unsigned int id = get_global_id(0); - unsigned int seed[4]; - - if (init != 0) { - seed[0] = seeds[id * 4]; - seed[1] = 0x50a7f451; - seed[2] = 0x5365417e; - seed[3] = 0xc3a4171a; - } - else { - seed[0] = seeds[id * 4]; - seed[1] = seeds[id * 4 + 1]; - seed[2] = seeds[id * 4 + 2]; - seed[3] = seeds[id * 4 + 3]; - } - - unsigned int numRandomNumbersPerItem = (numRandomNumbers + get_global_size(0) - 1) / get_global_size(0); - for (unsigned int i = 0; i < numRandomNumbersPerItem; i++) { - do - { - unsigned int alpha = (unsigned int)(seed[1] ^ (seed[1] << 11)); - seed[1] = seed[2]; - seed[2] = seed[3]; - seed[3] = seed[0]; - seed[0] = (seed[0] ^ (seed[0] >> 19)) ^ (alpha ^ (alpha >> 8)); - } while (seed[0] == ~0UL); - unsigned int pos = (get_group_id(0)*get_local_size(0)*numRandomNumbersPerItem) - + get_local_size(0) * i + get_local_id(0); - - if (pos >= numRandomNumbers) - break; - randomNumbers[pos] = normalizeRand*seed[0]; - } - - /* save the seeds for the time*/ - seeds[id * 4] = seed[0]; - seeds[id * 4 + 1] = seed[1]; - seeds[id * 4 + 2] = seed[2]; - seeds[id * 4 + 3] = seed[3]; - } - ) - /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % @@ -2641,7 +2568,7 @@ STRINGIFY( ) ; - const char* accelerateKernels2 = + const char *accelerateKernels2 = STRINGIFY( @@ -2740,17 +2667,17 @@ STRINGIFY( float4 cp = (float4) 0; - __local float *p = inputImageCache + (cacheIndex*number_channels); - cp.x = *(p); + __local CLQuantum *p = inputImageCache + (cacheIndex*number_channels); + cp.x = (float) *(p); if (number_channels > 2) { - cp.y = *(p + 1); - cp.z = *(p + 2); + cp.y = (float) *(p + 1); + cp.z = (float) *(p + 2); } if ((number_channels == 4) || (number_channels == 2)) { - cp.w = *(p + number_channels - 1); + cp.w = (float) *(p + number_channels - 1); float alpha = weight * QuantumScale * cp.w; @@ -2793,17 +2720,19 @@ STRINGIFY( if (itemID < actualNumPixelInThisChunk) { float density = densityCache[itemID]; + float gamma = gammaCache[itemID]; float4 filteredPixel = outputPixelCache[itemID]; if ((density != 0.0f) && (density != 1.0f)) { density = PerceptibleReciprocal(density); filteredPixel *= (float4) density; + gamma *= density; } if ((number_channels == 4) || (number_channels == 2)) { - float gamma = PerceptibleReciprocal(gammaCache[itemID]*density); + gamma = PerceptibleReciprocal(gamma); filteredPixel.x *= gamma; filteredPixel.y *= gamma; filteredPixel.z *= gamma; @@ -2899,17 +2828,17 @@ STRINGIFY( float4 cp = (float4)0.0f; - __local float *p = inputImageCache + cacheIndex; - cp.x = *(p); + __local CLQuantum *p = inputImageCache + cacheIndex; + cp.x = (float) *(p); if (number_channels > 2) { - cp.y = *(p + rangeLength); - cp.z = *(p + (rangeLength * 2)); + cp.y = (float) *(p + rangeLength); + cp.z = (float) *(p + (rangeLength * 2)); } if ((number_channels == 4) || (number_channels == 2)) { - cp.w = *(p + (rangeLength * (number_channels - 1))); + cp.w = (float) *(p + (rangeLength * (number_channels - 1))); float alpha = weight * QuantumScale * cp.w; @@ -2952,17 +2881,19 @@ STRINGIFY( if (itemID < actualNumPixelInThisChunk) { float density = densityCache[itemID]; + float gamma = gammaCache[itemID]; float4 filteredPixel = outputPixelCache[itemID]; if ((density != 0.0f) && (density != 1.0f)) { density = PerceptibleReciprocal(density); filteredPixel *= (float4) density; + gamma *= density; } if ((number_channels == 4) || (number_channels == 2)) { - float gamma = PerceptibleReciprocal(gammaCache[itemID]*density); + gamma = PerceptibleReciprocal(gamma); filteredPixel.x *= gamma; filteredPixel.y *= gamma; filteredPixel.z *= gamma; @@ -3138,8 +3069,7 @@ STRINGIFY( __kernel void UnsharpMask(const __global CLQuantum *image,const unsigned int number_channels, const ChannelType channel,__constant float *filter,const unsigned int width, const unsigned int columns,const unsigned int rows,__local float4 *pixels, - const float gain,const float threshold, const unsigned int justBlur, - __global CLQuantum *filteredImage) + const float gain,const float threshold,__global CLQuantum *filteredImage) { const unsigned int x = get_global_id(0); const unsigned int y = get_global_id(1); @@ -3199,15 +3129,13 @@ STRINGIFY( ++i; } - if (justBlur == 0) { // apply sharpening - float4 srcPixel = ReadFloat4(image, number_channels, columns, x, y, channel); - float4 diff = srcPixel - value; + float4 srcPixel = ReadFloat4(image, number_channels, columns, x, y, channel); + float4 diff = srcPixel - value; - float quantumThreshold = QuantumRange*threshold; + float quantumThreshold = QuantumRange*threshold; - int4 mask = isless(fabs(2.0f * diff), (float4)quantumThreshold); - value = select(srcPixel + diff * gain, srcPixel, mask); - } + int4 mask = isless(fabs(2.0f * diff), (float4)quantumThreshold); + value = select(srcPixel + diff * gain, srcPixel, mask); if ((x < columns) && (y < rows)) WriteFloat4(filteredImage, number_channels, columns, x, y, channel, value); diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index be9d46b8d6..f9c9a50f51 100644 --- a/MagickCore/accelerate.c +++ b/MagickCore/accelerate.c @@ -79,10 +79,6 @@ Include declarations. #include "MagickCore/string-private.h" #include "MagickCore/token.h" -#ifdef MAGICKCORE_CLPERFMARKER -#include "CLPerfMarker.h" -#endif - #define MAGICK_MAX(x,y) (((x) >= (y))?(x):(y)) #define MAGICK_MIN(x,y) (((x) <= (y))?(x):(y)) @@ -109,13 +105,6 @@ static const ResizeWeightingFunctionType supportedResizeWeighting[] = LastWeightingFunction }; -/* - Forward declarations. -*/ -static Image *ComputeUnsharpMaskImageSingle(const Image *image, - const double radius,const double sigma,const double gain, - const double threshold,int blurOnly,ExceptionInfo *exception); - /* Helper functions. */ @@ -218,35 +207,22 @@ static MagickBooleanType checkHistogramCondition(const Image *image, return(checkPixelIntensity(image,method)); } -static MagickBooleanType checkOpenCLEnvironment(ExceptionInfo* exception) +static MagickCLEnv getOpenCLEnvironment(ExceptionInfo* exception) { - MagickBooleanType - flag; - MagickCLEnv clEnv; - clEnv=GetDefaultOpenCLEnv(); + clEnv=GetCurrentOpenCLEnv(); + if (clEnv == (MagickCLEnv) NULL) + return((MagickCLEnv) NULL); - GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED, - sizeof(MagickBooleanType),&flag,exception); - if (flag != MagickFalse) - return(MagickFalse); + if (clEnv->enabled == MagickFalse) + return((MagickCLEnv) NULL); - GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED, - sizeof(MagickBooleanType),&flag,exception); - if (flag == MagickFalse) - { - if (InitOpenCLEnv(clEnv,exception) == MagickFalse) - return(MagickFalse); + if (InitializeOpenCL(clEnv,exception) == MagickFalse) + return((MagickCLEnv) NULL); - GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED, - sizeof(MagickBooleanType),&flag,exception); - if (flag != MagickFalse) - return(MagickFalse); - } - - return(MagickTrue); + return(clEnv); } /* pad the global workgroup size to the next multiple of @@ -258,7 +234,7 @@ inline static unsigned int padGlobalWorkgroupSizeToLocalWorkgroupSize( } static cl_mem createBuffer(const Image *image,CacheView *image_view, - MagickCLEnv clEnv,cl_context context,cl_mem_flags flags,void *pixels, + MagickCLEnv clEnv,MagickCLDevice device,cl_mem_flags flags,void *pixels, ExceptionInfo *exception) { cl_mem @@ -276,11 +252,11 @@ static cl_mem createBuffer(const Image *image,CacheView *image_view, void *hostPtr; - pixels=(void *) GetCacheViewVirtualPixels(image_view,0,0,image->columns, + pixels=(void *) GetCacheViewAuthenticPixels(image_view,0,0,image->columns, image->rows,exception); if (pixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(), + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); return (cl_mem) NULL; } @@ -295,11 +271,11 @@ static cl_mem createBuffer(const Image *image,CacheView *image_view, hostPtr=NULL; length=image->columns*image->rows*image->number_channels; - buffer=clEnv->library->clCreateBuffer(context,mem_flags,length* + buffer=clEnv->library->clCreateBuffer(device->context,mem_flags,length* sizeof(CLQuantum),hostPtr,&status); if (status != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(), + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), ResourceLimitWarning,"clEnv->library->clCreateBuffer failed.","."); } @@ -307,34 +283,34 @@ static cl_mem createBuffer(const Image *image,CacheView *image_view, } static inline cl_mem createReadBuffer(const Image *image,CacheView *image_view, - MagickCLEnv clEnv,cl_context context,ExceptionInfo *exception) + MagickCLEnv clEnv,MagickCLDevice device,ExceptionInfo *exception) { void *pixels; pixels=(void *) NULL; - return(createBuffer(image,image_view,clEnv,context,CL_MEM_READ_ONLY, + return(createBuffer(image,image_view,clEnv,device,CL_MEM_READ_ONLY, pixels,exception)); } static inline cl_mem createReadWriteBuffer(const Image *image, - CacheView *image_view,MagickCLEnv clEnv,cl_context context,void *pixels, + CacheView *image_view,MagickCLEnv clEnv,MagickCLDevice device,void *pixels, ExceptionInfo *exception) { - return(createBuffer(image,image_view,clEnv,context,CL_MEM_READ_WRITE,pixels, + return(createBuffer(image,image_view,clEnv,device,CL_MEM_READ_WRITE,pixels, exception)); } static inline cl_mem createWriteBuffer(Image *image,CacheView *image_view, - MagickCLEnv clEnv,cl_context context,void *pixels,ExceptionInfo *exception) + MagickCLEnv clEnv,MagickCLDevice device,void *pixels,ExceptionInfo *exception) { - return(createBuffer(image,image_view,clEnv,context,CL_MEM_WRITE_ONLY,pixels, + return(createBuffer(image,image_view,clEnv,device,CL_MEM_WRITE_ONLY,pixels, exception)); } static inline MagickBooleanType copyWriteBuffer(const Image *image, - MagickCLEnv clEnv,cl_command_queue queue,cl_mem buffer,void *pixels, - ExceptionInfo *exception) + MagickCLEnv clEnv,MagickCLDevice device,cl_command_queue queue,cl_mem buffer, + void *pixels,ExceptionInfo *exception) { cl_int status; @@ -342,16 +318,16 @@ static inline MagickBooleanType copyWriteBuffer(const Image *image, size_t length; - length=image->columns*image->rows*image->number_channels; + length=image->columns*image->rows*image->number_channels*sizeof(CLQuantum); if (ALIGNED(pixels,CLQuantum)) clEnv->library->clEnqueueMapBuffer(queue,buffer,CL_TRUE,CL_MAP_READ | - CL_MAP_WRITE,0,length*sizeof(CLQuantum),0,NULL,NULL,&status); + CL_MAP_WRITE,0,length,0,NULL,NULL,&status); else - status=clEnv->library->clEnqueueReadBuffer(queue,buffer,CL_TRUE,0,length* - sizeof(CLQuantum),pixels,0,NULL,NULL); + status=clEnv->library->clEnqueueReadBuffer(queue,buffer,CL_TRUE,0,length, + pixels,0,NULL,NULL); if (status != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(), + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), ResourceLimitWarning,"Reading output image from CL buffer failed.", "'%s'","."); return(MagickFalse); @@ -359,7 +335,7 @@ static inline MagickBooleanType copyWriteBuffer(const Image *image, return(MagickTrue); } -static cl_mem createKernelInfo(MagickCLEnv clEnv,cl_context context, +static cl_mem createKernelInfo(MagickCLEnv clEnv,MagickCLDevice device, cl_command_queue queue,const double radius,const double sigma,cl_uint *width, ExceptionInfo *exception) { @@ -386,17 +362,17 @@ static cl_mem createKernelInfo(MagickCLEnv clEnv,cl_context context, kernel=AcquireKernelInfo(geometry,exception); if (kernel == (KernelInfo *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(), + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), ResourceLimitWarning,"AcquireKernelInfo failed.","."); return((cl_mem) NULL); } - imageKernelBuffer=clEnv->library->clCreateBuffer(context,CL_MEM_READ_ONLY, - kernel->width*sizeof(float),NULL,&status); + imageKernelBuffer=clEnv->library->clCreateBuffer(device->context, + CL_MEM_READ_ONLY,kernel->width*sizeof(float),NULL,&status); if (status != CL_SUCCESS) { kernel=DestroyKernelInfo(kernel); - (void) OpenCLThrowMagickException(exception,GetMagickModule(), + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), ResourceLimitWarning,"clEnv->library->clCreateBuffer failed.","."); return((cl_mem) NULL); } @@ -408,14 +384,12 @@ static cl_mem createKernelInfo(MagickCLEnv clEnv,cl_context context, { kernel=DestroyKernelInfo(kernel); clEnv->library->clReleaseMemObject(imageKernelBuffer); - (void) OpenCLThrowMagickException(exception,GetMagickModule(), + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), ResourceLimitWarning,"clEnv->library->clEnqueueMapBuffer failed.","."); return((cl_mem) NULL); } for (i = 0; i < kernel->width; i++) - { kernelBufferPtr[i]=(float)kernel->values[i]; - } *width=(cl_uint) kernel->width; kernel=DestroyKernelInfo(kernel); @@ -425,7 +399,7 @@ static cl_mem createKernelInfo(MagickCLEnv clEnv,cl_context context, if (status != CL_SUCCESS) { clEnv->library->clReleaseMemObject(imageKernelBuffer); - (void) OpenCLThrowMagickException(exception,GetMagickModule(), + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), ResourceLimitWarning,"clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'","."); return((cl_mem) NULL); @@ -433,6 +407,83 @@ static cl_mem createKernelInfo(MagickCLEnv clEnv,cl_context context, return(imageKernelBuffer); } +static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv, + MagickCLDevice device,cl_command_queue queue,cl_mem imageBuffer, + cl_mem histogramBuffer,Image *image,const ChannelType channel, + ExceptionInfo *exception) +{ + MagickBooleanType + outputReady; + + cl_int + clStatus; + + cl_kernel + histogramKernel; + + cl_event + event; + + cl_uint + colorspace, + method; + + register ssize_t + i; + + size_t + global_work_size[2]; + + histogramKernel = NULL; + + outputReady = MagickFalse; + colorspace = image->colorspace; + method = image->intensity; + + /* get the OpenCL kernel */ + histogramKernel = AcquireOpenCLKernel(device,"Histogram"); + if (histogramKernel == NULL) + { + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + goto cleanup; + } + + /* set the kernel arguments */ + i = 0; + clStatus=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); + clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(ChannelType),&channel); + clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_uint),&colorspace); + clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_uint),&method); + clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&histogramBuffer); + if (clStatus != CL_SUCCESS) + { + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + goto cleanup; + } + + /* launch the kernel */ + global_work_size[0] = image->columns; + global_work_size[1] = image->rows; + + clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); + + if (clStatus != CL_SUCCESS) + { + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + goto cleanup; + } + RecordProfileData(device,HistogramKernel,event); + + outputReady = MagickTrue; + +cleanup: + + if (histogramKernel!=NULL) + RelinquishOpenCLKernel(histogramKernel); + + return(outputReady); +} + /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % @@ -445,7 +496,7 @@ static cl_mem createKernelInfo(MagickCLEnv clEnv,cl_context context, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static Image *ComputeAddNoiseImage(const Image *image, +static Image *ComputeAddNoiseImage(const Image *image,MagickCLEnv clEnv, const NoiseType noise_type,ExceptionInfo *exception) { CacheView @@ -455,9 +506,6 @@ static Image *ComputeAddNoiseImage(const Image *image, cl_command_queue queue; - cl_context - context; - cl_float attenuate; @@ -489,8 +537,8 @@ static Image *ComputeAddNoiseImage(const Image *image, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; Image *filteredImage; @@ -512,12 +560,11 @@ static Image *ComputeAddNoiseImage(const Image *image, filteredImageBuffer = NULL; addNoiseKernel = NULL; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); - image_view=AcquireVirtualCacheView(image,exception); - imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); + image_view=AcquireAuthenticCacheView(image,exception); + imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -526,13 +573,13 @@ static Image *ComputeAddNoiseImage(const Image *image, goto cleanup; if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, - context,filteredPixels,exception); + device,filteredPixels,exception); if (filteredImageBuffer == (void *) NULL) goto cleanup; @@ -566,18 +613,16 @@ static Image *ComputeAddNoiseImage(const Image *image, numRandomNumberPerPixel+=numRandPerChannel; } - addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoise"); + addNoiseKernel = AcquireOpenCLKernel(device,"AddNoise"); if (addNoiseKernel == NULL) { - (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } { - cl_uint computeUnitCount; cl_uint workItemCount; - clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &computeUnitCount, NULL); - workItemCount = computeUnitCount * 2 * 256; // 256 work items per group, 2 groups per CU + workItemCount = device->max_compute_units * 2 * 256; // 256 work items per group, 2 groups per CU inputPixelCount = (cl_int) (image->columns * image->rows); pixelsPerWorkitem = (inputPixelCount + workItemCount - 1) / workItemCount; pixelsPerWorkitem = ((pixelsPerWorkitem + 3) / 4) * 4; @@ -615,33 +660,31 @@ static Image *ComputeAddNoiseImage(const Image *image, clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } clStatus=clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,1,NULL,global_work_size,local_work_size,0,NULL,&event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,AddNoiseKernel,event); - clEnv->library->clReleaseEvent(event); - if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) + RecordProfileData(device,AddNoiseKernel,event); + if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) goto cleanup; outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (filteredImage_view != NULL) filteredImage_view=DestroyCacheView(filteredImage_view); - if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue); - if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel); + if (queue!=NULL) RelinquishOpenCLCommandQueue(device, queue); + if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(addNoiseKernel); if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); if (outputReady == MagickFalse && filteredImage != NULL) @@ -656,14 +699,20 @@ MagickExport Image *AccelerateAddNoiseImage(const Image *image, Image *filteredImage; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateCondition(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) + if (checkAccelerateCondition(image) == MagickFalse) return((Image *) NULL); - filteredImage=ComputeAddNoiseImage(image,noise_type,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return((Image *) NULL); + + filteredImage=ComputeAddNoiseImage(image,clEnv,noise_type,exception); return(filteredImage); } @@ -679,8 +728,8 @@ MagickExport Image *AccelerateAddNoiseImage(const Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static Image *ComputeBlurImage(const Image* image,const double radius, - const double sigma,ExceptionInfo *exception) +static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv, + const double radius,const double sigma,ExceptionInfo *exception) { CacheView *filteredImage_view, @@ -689,9 +738,6 @@ static Image *ComputeBlurImage(const Image* image,const double radius, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -720,8 +766,8 @@ static Image *ComputeBlurImage(const Image* image,const double radius, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -732,7 +778,6 @@ static Image *ComputeBlurImage(const Image* image,const double radius, void *filteredPixels; - context = NULL; filteredImage = NULL; filteredImage_view = NULL; imageBuffer = NULL; @@ -746,12 +791,11 @@ static Image *ComputeBlurImage(const Image* image,const double radius, outputReady = MagickFalse; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); - image_view=AcquireVirtualCacheView(image,exception); - imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); + image_view=AcquireAuthenticCacheView(image,exception); + imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -760,44 +804,44 @@ static Image *ComputeBlurImage(const Image* image,const double radius, goto cleanup; if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, - context,filteredPixels,exception); + device,filteredPixels,exception); if (filteredImageBuffer == (void *) NULL) goto cleanup; - imageKernelBuffer=createKernelInfo(clEnv,context,queue,radius,sigma, + imageKernelBuffer=createKernelInfo(clEnv,device,queue,radius,sigma, &kernelWidth,exception); { /* create temp buffer */ { length = image->columns * image->rows; - tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(cl_float4), NULL, &clStatus); + tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length * sizeof(cl_float4), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } } /* get the OpenCL kernels */ { - blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow"); + blurRowKernel = AcquireOpenCLKernel(device,"BlurRow"); if (blurRowKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; - blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumn"); + blurColumnKernel = AcquireOpenCLKernel(device,"BlurColumn"); if (blurColumnKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; } @@ -824,7 +868,7 @@ static Image *ComputeBlurImage(const Image* image,const double radius, clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -842,12 +886,10 @@ static Image *ComputeBlurImage(const Image* image,const double radius, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,BlurRowKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,BlurRowKernel,event); } } @@ -869,7 +911,7 @@ static Image *ComputeBlurImage(const Image* image,const double radius, clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -887,25 +929,22 @@ static Image *ComputeBlurImage(const Image* image,const double radius, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,BlurColumnKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,BlurColumnKernel,event); } } } /* get result */ - if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) + if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) goto cleanup; outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (filteredImage_view != NULL) @@ -915,37 +954,34 @@ cleanup: if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer); if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); - if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel); - if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (blurRowKernel!=NULL) RelinquishOpenCLKernel(blurRowKernel); + if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(blurColumnKernel); + if (queue != NULL) RelinquishOpenCLCommandQueue(device, queue); if (outputReady == MagickFalse && filteredImage != NULL) filteredImage=DestroyImage(filteredImage); return(filteredImage); } -static Image* ComputeBlurImageSingle(const Image* image, - const double radius,const double sigma,ExceptionInfo *exception) -{ - return ComputeUnsharpMaskImageSingle(image,radius,sigma,0.0,0.0,1,exception); -} - MagickExport Image* AccelerateBlurImage(const Image *image, const double radius,const double sigma,ExceptionInfo *exception) { Image *filteredImage; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateCondition(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return NULL; + if (checkAccelerateCondition(image) == MagickFalse) + return((Image *) NULL); - if (radius < 12.1) - filteredImage=ComputeBlurImageSingle(image,radius,sigma,exception); - else - filteredImage=ComputeBlurImage(image,radius,sigma,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return((Image *) NULL); + + filteredImage=ComputeBlurImage(image,clEnv,radius,sigma,exception); return(filteredImage); } @@ -962,12 +998,12 @@ MagickExport Image* AccelerateBlurImage(const Image *image, */ static MagickBooleanType LaunchCompositeKernel(MagickCLEnv clEnv, - cl_command_queue queue,cl_mem imageBuffer,const unsigned int inputWidth, - const unsigned int inputHeight,const unsigned int matte, - const ChannelType channel,const CompositeOperator compose, - const cl_mem compositeImageBuffer,const unsigned int compositeWidth, - const unsigned int compositeHeight,const float destination_dissolve, - const float source_dissolve,ExceptionInfo *magick_unused(exception)) + MagickCLDevice device,cl_command_queue queue,cl_mem imageBuffer, + const unsigned int inputWidth,const unsigned int inputHeight, + const unsigned int matte,const ChannelType channel, + const CompositeOperator compose,const cl_mem compositeImageBuffer, + const unsigned int compositeWidth,const unsigned int compositeHeight, + const float destination_dissolve,const float source_dissolve) { cl_int clStatus; @@ -988,10 +1024,7 @@ static MagickBooleanType LaunchCompositeKernel(MagickCLEnv clEnv, unsigned int composeOp; - magick_unreferenced(exception); - - compositeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, - "Composite"); + compositeKernel = AcquireOpenCLKernel(device,"Composite"); k = 0; clStatus=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(cl_mem),(void*)&imageBuffer); @@ -1017,17 +1050,16 @@ static MagickBooleanType LaunchCompositeKernel(MagickCLEnv clEnv, (unsigned int) local_work_size[0]); global_work_size[1] = inputHeight; clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, compositeKernel, 2, NULL, - global_work_size, local_work_size, 0, NULL, &event); + global_work_size, local_work_size, 0, NULL, &event); - RecordProfileData(clEnv,CompositeKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,CompositeKernel,event); - RelinquishOpenCLKernel(clEnv, compositeKernel); + RelinquishOpenCLKernel(compositeKernel); return((clStatus==CL_SUCCESS) ? MagickTrue : MagickFalse); } -static MagickBooleanType ComputeCompositeImage(Image *image, +static MagickBooleanType ComputeCompositeImage(Image *image,MagickCLEnv clEnv, const CompositeOperator compose,const Image *compositeImage, const float destination_dissolve,const float source_dissolve,ExceptionInfo *exception) { @@ -1037,9 +1069,6 @@ static MagickBooleanType ComputeCompositeImage(Image *image, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -1057,8 +1086,8 @@ static MagickBooleanType ComputeCompositeImage(Image *image, outputReady, status; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -1072,16 +1101,15 @@ static MagickBooleanType ComputeCompositeImage(Image *image, imageBuffer = NULL; compositeImageBuffer = NULL; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); /* Create and initialize OpenCL buffers. */ image_view=AcquireAuthenticCacheView(image,exception); inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); if (inputPixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",image->filename); goto cleanup; } @@ -1099,11 +1127,11 @@ static MagickBooleanType ComputeCompositeImage(Image *image, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, + imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } @@ -1113,7 +1141,7 @@ static MagickBooleanType ComputeCompositeImage(Image *image, composePixels = AcquirePixelCachePixels(compositeImage, &length, exception); if (composePixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",compositeImage->filename); goto cleanup; } @@ -1131,24 +1159,23 @@ static MagickBooleanType ComputeCompositeImage(Image *image, } /* create a CL buffer from image pixel buffer */ length = compositeImage->columns * compositeImage->rows; - compositeImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, + compositeImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)composePixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - status = LaunchCompositeKernel(clEnv,queue,imageBuffer, + status = LaunchCompositeKernel(clEnv,device,queue,imageBuffer, (unsigned int) image->columns, (unsigned int) image->rows, (unsigned int) (image->alpha_trait > CopyPixelTrait) ? 1 : 0, image->channel_mask, compose, compositeImageBuffer, (unsigned int) compositeImage->columns, (unsigned int) compositeImage->rows, - destination_dissolve,source_dissolve, - exception); + destination_dissolve,source_dissolve); if (status==MagickFalse) goto cleanup; @@ -1173,7 +1200,7 @@ cleanup: image_view=DestroyCacheView(image_view); if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (compositeImageBuffer!=NULL) clEnv->library->clReleaseMemObject(compositeImageBuffer); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (queue != NULL) RelinquishOpenCLCommandQueue(device,queue); return(outputReady); } @@ -1186,17 +1213,16 @@ MagickExport MagickBooleanType AccelerateCompositeImage(Image *image, MagickBooleanType status; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateConditionRGBA(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return(MagickFalse); - /* only support images with the size for now */ if ((image->columns != composite->columns) || (image->rows != composite->rows)) - return MagickFalse; + return(MagickFalse); switch(compose) { @@ -1205,11 +1231,18 @@ MagickExport MagickBooleanType AccelerateCompositeImage(Image *image, break; default: // unsupported compose operator, quit - return MagickFalse; + return(MagickFalse); }; - status=ComputeCompositeImage(image,compose,composite,destination_dissolve, - source_dissolve,exception); + if (checkAccelerateConditionRGBA(image) == MagickFalse) + return(MagickFalse); + + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return(MagickFalse); + + status=ComputeCompositeImage(image,clEnv,compose,composite, + destination_dissolve,source_dissolve,exception); return(status); } @@ -1225,7 +1258,7 @@ MagickExport MagickBooleanType AccelerateCompositeImage(Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static MagickBooleanType ComputeContrastImage(Image *image, +static MagickBooleanType ComputeContrastImage(Image *image,MagickCLEnv clEnv, const MagickBooleanType sharpen,ExceptionInfo *exception) { CacheView @@ -1234,9 +1267,6 @@ static MagickBooleanType ComputeContrastImage(Image *image, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -1255,8 +1285,8 @@ static MagickBooleanType ComputeContrastImage(Image *image, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -1272,22 +1302,19 @@ static MagickBooleanType ComputeContrastImage(Image *image, *inputPixels; outputReady = MagickFalse; - clEnv = NULL; inputPixels = NULL; - context = NULL; imageBuffer = NULL; filterKernel = NULL; queue = NULL; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); + device = GetOpenCLDevice(clEnv); /* Create and initialize OpenCL buffers. */ image_view=AcquireAuthenticCacheView(image,exception); inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); if (inputPixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); goto cleanup; } @@ -1304,17 +1331,17 @@ static MagickBooleanType ComputeContrastImage(Image *image, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); + imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - filterKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Contrast"); + filterKernel = AcquireOpenCLKernel(device,"Contrast"); if (filterKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -1325,23 +1352,21 @@ static MagickBooleanType ComputeContrastImage(Image *image, clStatus|=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_uint),&uSharpen); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } global_work_size[0] = image->columns; global_work_size[1] = image->rows; /* launch the kernel */ - queue = AcquireOpenCLCommandQueue(clEnv); + queue = AcquireOpenCLCommandQueue(device); clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, filterKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,ContrastKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,ContrastKernel,event); if (ALIGNED(inputPixels,CLPixelPacket)) { @@ -1355,19 +1380,18 @@ static MagickBooleanType ComputeContrastImage(Image *image, } if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady=SyncCacheViewAuthenticPixels(image_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); - if (filterKernel!=NULL) RelinquishOpenCLKernel(clEnv, filterKernel); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (filterKernel!=NULL) RelinquishOpenCLKernel(filterKernel); + if (queue != NULL) RelinquishOpenCLCommandQueue(device,queue); return(outputReady); } @@ -1377,14 +1401,20 @@ MagickExport MagickBooleanType AccelerateContrastImage(Image *image, MagickBooleanType status; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateConditionRGBA(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) + if (checkAccelerateConditionRGBA(image) == MagickFalse) return(MagickFalse); - status=ComputeContrastImage(image,sharpen,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return(MagickFalse); + + status=ComputeContrastImage(image,clEnv,sharpen,exception); return(status); } @@ -1400,87 +1430,9 @@ MagickExport MagickBooleanType AccelerateContrastImage(Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv, - cl_command_queue queue,cl_mem imageBuffer,cl_mem histogramBuffer, - Image *image,const ChannelType channel,ExceptionInfo *exception) -{ - MagickBooleanType - outputReady; - - cl_int - clStatus; - - cl_kernel - histogramKernel; - - cl_event - event; - - cl_uint - colorspace, - method; - - register ssize_t - i; - - size_t - global_work_size[2]; - - histogramKernel = NULL; - - outputReady = MagickFalse; - colorspace = image->colorspace; - method = image->intensity; - - /* get the OpenCL kernel */ - histogramKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Histogram"); - if (histogramKernel == NULL) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); - goto cleanup; - } - - /* set the kernel arguments */ - i = 0; - clStatus=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); - clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(ChannelType),&channel); - clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_uint),&colorspace); - clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_uint),&method); - clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&histogramBuffer); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); - goto cleanup; - } - - /* launch the kernel */ - global_work_size[0] = image->columns; - global_work_size[1] = image->rows; - - clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); - - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); - goto cleanup; - } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,HistogramKernel,event); - clEnv->library->clReleaseEvent(event); - - outputReady = MagickTrue; - -cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); - - if (histogramKernel!=NULL) - RelinquishOpenCLKernel(clEnv, histogramKernel); - - return(outputReady); -} - static MagickBooleanType ComputeContrastStretchImage(Image *image, - const double black_point,const double white_point,ExceptionInfo *exception) + MagickCLEnv clEnv,const double black_point,const double white_point, + ExceptionInfo *exception) { #define ContrastStretchImageTag "ContrastStretch/Image" #define MaxRange(color) ((MagickRealType) ScaleQuantumToMap((Quantum) (color))) @@ -1491,9 +1443,6 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -1526,8 +1475,8 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, outputReady, status; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -1553,7 +1502,6 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, stretchMapBuffer = NULL; histogramKernel = NULL; stretchKernel = NULL; - context = NULL; queue = NULL; outputReady = MagickFalse; @@ -1568,9 +1516,8 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, /* * initialize opencl env */ - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); /* Allocate and initialize histogram arrays. @@ -1602,7 +1549,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, if (inputPixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); goto cleanup; } /* If the host pointer is aligned to the size of CLPixelPacket, @@ -1618,10 +1565,10 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); + imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } @@ -1640,14 +1587,14 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, } /* create a CL buffer for histogram */ length = (MaxMap+1); - histogramBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus); + histogramBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, image->channel_mask, exception); + status = LaunchHistogramKernel(clEnv, device, queue, imageBuffer, histogramBuffer, image, image->channel_mask,exception); if (status == MagickFalse) goto cleanup; @@ -1664,7 +1611,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, } if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } @@ -1674,7 +1621,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } } @@ -1943,7 +1890,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } #endif @@ -1961,18 +1908,18 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, } /* create a CL buffer for stretch_map */ length = (MaxMap+1); - stretchMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus); + stretchMapBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } /* get the OpenCL kernel */ - stretchKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ContrastStretch"); + stretchKernel = AcquireOpenCLKernel(device,"ContrastStretch"); if (stretchKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -1985,7 +1932,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&black); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -1997,13 +1944,10 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - - RecordProfileData(clEnv,ContrastStretchKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,ContrastStretchKernel,event); /* read the data back */ if (ALIGNED(inputPixels,CLPixelPacket)) @@ -2018,14 +1962,13 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image, } if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady=SyncCacheViewAuthenticPixels(image_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); @@ -2045,12 +1988,12 @@ cleanup: if (histogramKernel!=NULL) - RelinquishOpenCLKernel(clEnv, histogramKernel); + RelinquishOpenCLKernel(histogramKernel); if (stretchKernel!=NULL) - RelinquishOpenCLKernel(clEnv, stretchKernel); + RelinquishOpenCLKernel(stretchKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(clEnv, queue); + RelinquishOpenCLCommandQueue(device, queue); return(outputReady); } @@ -2062,15 +2005,22 @@ MagickExport MagickBooleanType AccelerateContrastStretchImage( MagickBooleanType status; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); if ((checkAccelerateConditionRGBA(image) == MagickFalse) || - (checkHistogramCondition(image,image->intensity) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) + (checkHistogramCondition(image,image->intensity) == MagickFalse)) return(MagickFalse); - status=ComputeContrastStretchImage(image,black_point,white_point,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return(MagickFalse); + + status=ComputeContrastStretchImage(image,clEnv,black_point,white_point, + exception); return(status); } @@ -2086,8 +2036,8 @@ MagickExport MagickBooleanType AccelerateContrastStretchImage( %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, - ExceptionInfo *exception) +static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv, + const KernelInfo *kernel,ExceptionInfo *exception) { CacheView *filteredImage_view, @@ -2096,15 +2046,12 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, cl_command_queue queue; - cl_context - context; + cl_event + event; cl_kernel clkernel; - cl_event - event; - cl_int clStatus; @@ -2116,9 +2063,6 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, cl_mem_flags mem_flags; - cl_ulong - deviceLocalMemorySize; - const void *inputPixels; @@ -2131,8 +2075,8 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -2158,7 +2102,6 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, *hostPtr; /* intialize all CL objects to NULL */ - context = NULL; imageBuffer = NULL; filteredImageBuffer = NULL; convolutionKernel = NULL; @@ -2168,15 +2111,14 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, filteredImage = NULL; filteredImage_view = NULL; outputReady = MagickFalse; - - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - image_view=AcquireVirtualCacheView(image,exception); - inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); + device = GetOpenCLDevice(clEnv); + + image_view=AcquireAuthenticCacheView(image,exception); + inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); if (inputPixels == (const void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); goto cleanup; } @@ -2195,10 +2137,10 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); + imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } @@ -2206,14 +2148,14 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); if (filteredPixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -2229,28 +2171,28 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); + filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } kernelSize = (unsigned int) (kernel->width * kernel->height); - convolutionKernel = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize * sizeof(float), NULL, &clStatus); + convolutionKernel = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - queue = AcquireOpenCLCommandQueue(clEnv); + queue = AcquireOpenCLCommandQueue(device); kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, convolutionKernel, CL_TRUE, CL_MAP_WRITE, 0, kernelSize * sizeof(float) , 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); goto cleanup; } for (i = 0; i < kernelSize; i++) @@ -2260,12 +2202,9 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - - deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv); /* Compute the local memory requirement for a 16x16 workgroup. If it's larger than 16k, reduce the workgroup size to 8x8 */ @@ -2274,20 +2213,20 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket) + kernel->width*kernel->height*sizeof(float); - if (localMemoryRequirement > deviceLocalMemorySize) + if (localMemoryRequirement > device->local_memory_size) { localGroupSize[0] = 8; localGroupSize[1] = 8; localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket) + kernel->width*kernel->height*sizeof(float); } - if (localMemoryRequirement <= deviceLocalMemorySize) + if (localMemoryRequirement <= device->local_memory_size) { /* get the OpenCL kernel */ - clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ConvolveOptimized"); + clkernel = AcquireOpenCLKernel(device,"ConvolveOptimized"); if (clkernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -2311,7 +2250,7 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, kernel->width*kernel->height*sizeof(float),NULL); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -2320,22 +2259,21 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, global_work_size[1] = ((image->rows + localGroupSize[1] - 1)/localGroupSize[1]) * localGroupSize[1]; /* launch the kernel */ - clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, &event); + clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,ConvolveKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,ConvolveOptimizedKernel,event); } else { /* get the OpenCL kernel */ - clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Convolve"); + clkernel = AcquireOpenCLKernel(device,"Convolve"); if (clkernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -2357,7 +2295,7 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&image->channel_mask); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -2369,13 +2307,11 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,ConvolveKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,ConvolveKernel,event); if (ALIGNED(filteredPixels,CLPixelPacket)) { @@ -2389,14 +2325,13 @@ static Image *ComputeConvolveImage(const Image* image,const KernelInfo *kernel, } if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (filteredImage_view != NULL) @@ -2412,10 +2347,10 @@ cleanup: clEnv->library->clReleaseMemObject(convolutionKernel); if (clkernel != NULL) - RelinquishOpenCLKernel(clEnv, clkernel); + RelinquishOpenCLKernel(clkernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(clEnv, queue); + RelinquishOpenCLCommandQueue(device, queue); if (outputReady == MagickFalse) { @@ -2465,7 +2400,7 @@ MagickExport Image *AccelerateConvolveImage(const Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static Image *ComputeDespeckleImage(const Image *image, +static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv, ExceptionInfo*exception) { static const int @@ -2479,9 +2414,6 @@ static Image *ComputeDespeckleImage(const Image *image, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -2513,8 +2445,8 @@ static Image *ComputeDespeckleImage(const Image *image, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -2531,27 +2463,25 @@ static Image *ComputeDespeckleImage(const Image *image, *hostPtr; outputReady = MagickFalse; - clEnv = NULL; inputPixels = NULL; filteredImage = NULL; filteredImage_view = NULL; filteredPixels = NULL; - context = NULL; imageBuffer = NULL; filteredImageBuffer = NULL; hullPass1 = NULL; hullPass2 = NULL; queue = NULL; tempImageBuffer[0] = tempImageBuffer[1] = NULL; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); - image_view=AcquireVirtualCacheView(image,exception); - inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); + image_view=AcquireAuthenticCacheView(image,exception); + inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); if (inputPixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); goto cleanup; } @@ -2565,10 +2495,10 @@ static Image *ComputeDespeckleImage(const Image *image, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); + imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } @@ -2576,10 +2506,10 @@ static Image *ComputeDespeckleImage(const Image *image, length = image->columns * image->rows; for (k = 0; k < 2; k++) { - tempImageBuffer[k] = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), NULL, &clStatus); + tempImageBuffer[k] = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } } @@ -2588,14 +2518,14 @@ static Image *ComputeDespeckleImage(const Image *image, assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); if (filteredPixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -2611,15 +2541,15 @@ static Image *ComputeDespeckleImage(const Image *image, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); + filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - hullPass1 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass1"); - hullPass2 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass2"); + hullPass1 = AcquireOpenCLKernel(device,"HullPass1"); + hullPass2 = AcquireOpenCLKernel(device,"HullPass2"); clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)&imageBuffer); clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,sizeof(cl_mem),(void *)(tempImageBuffer+1)); @@ -2631,7 +2561,7 @@ static Image *ComputeDespeckleImage(const Image *image, clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -2645,7 +2575,7 @@ static Image *ComputeDespeckleImage(const Image *image, clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -2669,28 +2599,26 @@ static Image *ComputeDespeckleImage(const Image *image, clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,HullPass1Kernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,HullPass1Kernel,event); /* launch the kernel */ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,HullPass2Kernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,HullPass2Kernel,event); if (k == 0) clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)(tempImageBuffer)); @@ -2703,28 +2631,26 @@ static Image *ComputeDespeckleImage(const Image *image, clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,HullPass1Kernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,HullPass1Kernel,event); /* launch the kernel */ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,HullPass2Kernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,HullPass2Kernel,event); offset.s[0] = -X[k]; offset.s[1] = -Y[k]; @@ -2735,28 +2661,26 @@ static Image *ComputeDespeckleImage(const Image *image, clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,HullPass1Kernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,HullPass1Kernel,event); /* launch the kernel */ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,HullPass2Kernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,HullPass2Kernel,event); offset.s[0] = X[k]; offset.s[1] = Y[k]; @@ -2771,28 +2695,26 @@ static Image *ComputeDespeckleImage(const Image *image, if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,HullPass1Kernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,HullPass1Kernel,event); /* launch the kernel */ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,HullPass2Kernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,HullPass2Kernel,event); } if (ALIGNED(filteredPixels,CLPixelPacket)) @@ -2807,28 +2729,27 @@ static Image *ComputeDespeckleImage(const Image *image, } if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (filteredImage_view != NULL) filteredImage_view=DestroyCacheView(filteredImage_view); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (queue != NULL) RelinquishOpenCLCommandQueue(device, queue); if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); for (k = 0; k < 2; k++) { if (tempImageBuffer[k]!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer[k]); } if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); - if (hullPass1!=NULL) RelinquishOpenCLKernel(clEnv, hullPass1); - if (hullPass2!=NULL) RelinquishOpenCLKernel(clEnv, hullPass2); + if (hullPass1!=NULL) RelinquishOpenCLKernel(hullPass1); + if (hullPass2!=NULL) RelinquishOpenCLKernel(hullPass2); if (outputReady == MagickFalse && filteredImage != NULL) filteredImage=DestroyImage(filteredImage); return(filteredImage); @@ -2840,14 +2761,20 @@ MagickExport Image *AccelerateDespeckleImage(const Image* image, Image *filteredImage; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateConditionRGBA(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return NULL; + if (checkAccelerateConditionRGBA(image) == MagickFalse) + return((Image *) NULL); - filteredImage=ComputeDespeckleImage(image,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return((Image *) NULL); + + filteredImage=ComputeDespeckleImage(image,clEnv,exception); return(filteredImage); } @@ -2863,7 +2790,7 @@ MagickExport Image *AccelerateDespeckleImage(const Image* image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static MagickBooleanType ComputeEqualizeImage(Image *image, +static MagickBooleanType ComputeEqualizeImage(Image *image,MagickCLEnv clEnv, ExceptionInfo *exception) { #define EqualizeImageTag "Equalize/Image" @@ -2874,9 +2801,6 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -2908,8 +2832,8 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, outputReady, status; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -2936,7 +2860,6 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, equalizeMapBuffer = NULL; histogramKernel = NULL; equalizeKernel = NULL; - context = NULL; queue = NULL; outputReady = MagickFalse; @@ -2948,9 +2871,8 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, /* * initialize opencl env */ - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); /* Allocate and initialize histogram arrays. @@ -2970,7 +2892,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, if (inputPixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); goto cleanup; } /* If the host pointer is aligned to the size of CLPixelPacket, @@ -2986,10 +2908,10 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); + imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } @@ -3008,14 +2930,14 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, } /* create a CL buffer for histogram */ length = (MaxMap+1); - histogramBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus); + histogramBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, image->channel_mask, exception); + status = LaunchHistogramKernel(clEnv, device, queue, imageBuffer, histogramBuffer, image, image->channel_mask, exception); if (status == MagickFalse) goto cleanup; @@ -3032,7 +2954,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, } if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } @@ -3042,7 +2964,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } } @@ -3188,7 +3110,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } #endif @@ -3206,18 +3128,18 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, } /* create a CL buffer for eqaulize_map */ length = (MaxMap+1); - equalizeMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus); + equalizeMapBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } /* get the OpenCL kernel */ - equalizeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Equalize"); + equalizeKernel = AcquireOpenCLKernel(device,"Equalize"); if (equalizeKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -3230,7 +3152,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&black); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -3242,12 +3164,10 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,EqualizeKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,EqualizeKernel,event); /* read the data back */ if (ALIGNED(inputPixels,CLPixelPacket)) @@ -3262,14 +3182,13 @@ static MagickBooleanType ComputeEqualizeImage(Image *image, } if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady=SyncCacheViewAuthenticPixels(image_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); @@ -3290,12 +3209,12 @@ cleanup: histogram=(cl_uint4 *) RelinquishMagickMemory(histogram); if (histogramKernel!=NULL) - RelinquishOpenCLKernel(clEnv, histogramKernel); + RelinquishOpenCLKernel(histogramKernel); if (equalizeKernel!=NULL) - RelinquishOpenCLKernel(clEnv, equalizeKernel); + RelinquishOpenCLKernel(equalizeKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(clEnv, queue); + RelinquishOpenCLCommandQueue(device, queue); return(outputReady); } @@ -3306,15 +3225,21 @@ MagickExport MagickBooleanType AccelerateEqualizeImage(Image *image, MagickBooleanType status; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); if ((checkAccelerateConditionRGBA(image) == MagickFalse) || - (checkHistogramCondition(image,image->intensity) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) + (checkHistogramCondition(image,image->intensity) == MagickFalse)) return(MagickFalse); - status=ComputeEqualizeImage(image,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return(MagickFalse); + + status=ComputeEqualizeImage(image,clEnv,exception); return(status); } @@ -3330,7 +3255,7 @@ MagickExport MagickBooleanType AccelerateEqualizeImage(Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static MagickBooleanType ComputeFunctionImage(Image *image, +static MagickBooleanType ComputeFunctionImage(Image *image,MagickCLEnv clEnv, const MagickFunction function,const size_t number_parameters, const double *parameters,ExceptionInfo *exception) { @@ -3340,9 +3265,6 @@ static MagickBooleanType ComputeFunctionImage(Image *image, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -3365,8 +3287,8 @@ static MagickBooleanType ComputeFunctionImage(Image *image, MagickBooleanType status; - MagickCLEnv - clEnv; + MagickCLDevice + device; size_t globalWorkSize[2]; @@ -3379,36 +3301,34 @@ static MagickBooleanType ComputeFunctionImage(Image *image, status = MagickFalse; - context = NULL; clkernel = NULL; queue = NULL; imageBuffer = NULL; parametersBuffer = NULL; pixels = NULL; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); + device = GetOpenCLDevice(clEnv); image_view=AcquireAuthenticCacheView(image,exception); - imageBuffer=createReadWriteBuffer(image,image_view,clEnv,context,pixels, + imageBuffer=createReadWriteBuffer(image,image_view,clEnv,device,pixels, exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; - parametersBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, number_parameters * sizeof(float), NULL, &clStatus); + parametersBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, number_parameters * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - queue = AcquireOpenCLCommandQueue(clEnv); + queue = AcquireOpenCLCommandQueue(device); parametersBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, parametersBuffer, CL_TRUE, CL_MAP_WRITE, 0, number_parameters * sizeof(float) , 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); goto cleanup; } for (i = 0; i < number_parameters; i++) @@ -3418,15 +3338,14 @@ static MagickBooleanType ComputeFunctionImage(Image *image, clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, parametersBuffer, parametersBufferPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ComputeFunction"); + clkernel = AcquireOpenCLKernel(device,"ComputeFunction"); if (clkernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -3442,7 +3361,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image, clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)¶metersBuffer); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -3452,25 +3371,22 @@ static MagickBooleanType ComputeFunctionImage(Image *image, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, globalWorkSize, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,ComputeFunctionKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,ComputeFunctionKernel,event); - if (copyWriteBuffer(image,clEnv,queue,imageBuffer,pixels,exception) == MagickFalse) + if (copyWriteBuffer(image,clEnv,device,queue,imageBuffer,pixels,exception) == MagickFalse) goto cleanup; status=SyncCacheViewAuthenticPixels(image_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); - if (clkernel != NULL) RelinquishOpenCLKernel(clEnv, clkernel); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (clkernel != NULL) RelinquishOpenCLKernel(clkernel); + if (queue != NULL) RelinquishOpenCLCommandQueue(device, queue); if (imageBuffer != NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (parametersBuffer != NULL) clEnv->library->clReleaseMemObject(parametersBuffer); @@ -3484,15 +3400,21 @@ MagickExport MagickBooleanType AccelerateFunctionImage(Image *image, MagickBooleanType status; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateCondition(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) + if (checkAccelerateCondition(image) == MagickFalse) return(MagickFalse); - status=ComputeFunctionImage(image,function,number_parameters,parameters, - exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return(MagickFalse); + + status=ComputeFunctionImage(image,clEnv,function,number_parameters, + parameters,exception); return(status); } @@ -3508,7 +3430,7 @@ MagickExport MagickBooleanType AccelerateFunctionImage(Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static MagickBooleanType ComputeGrayscaleImage(Image *image, +static MagickBooleanType ComputeGrayscaleImage(Image *image,MagickCLEnv clEnv, const PixelIntensityMethod method,ExceptionInfo *exception) { CacheView @@ -3517,9 +3439,6 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -3540,8 +3459,8 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; register ssize_t i; @@ -3561,24 +3480,23 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image, /* * initialize opencl env */ - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); /* Create and initialize OpenCL buffers. inputPixels = AcquirePixelCachePixels(image, &length, exception); assume this will get a writable image */ image_view=AcquireAuthenticCacheView(image,exception); - imageBuffer=createReadWriteBuffer(image,image_view,clEnv,context,inputPixels, + imageBuffer=createReadWriteBuffer(image,image_view,clEnv,device,inputPixels, exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; - grayscaleKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Grayscale"); + grayscaleKernel = AcquireOpenCLKernel(device,"Grayscale"); if (grayscaleKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -3593,7 +3511,7 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image, clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&intensityMethod); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -3605,30 +3523,27 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image, clStatus=clEnv->library->clEnqueueNDRangeKernel(queue, grayscaleKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,GrayScaleKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,GrayScaleKernel,event); } - if (copyWriteBuffer(image,clEnv,queue,imageBuffer,inputPixels,exception) == MagickFalse) + if (copyWriteBuffer(image,clEnv,device,queue,imageBuffer,inputPixels,exception) == MagickFalse) goto cleanup; outputReady=SyncCacheViewAuthenticPixels(image_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (grayscaleKernel!=NULL) - RelinquishOpenCLKernel(clEnv, grayscaleKernel); + RelinquishOpenCLKernel(grayscaleKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(clEnv, queue); + RelinquishOpenCLCommandQueue(device, queue); return( outputReady); } @@ -3639,12 +3554,14 @@ MagickExport MagickBooleanType AccelerateGrayscaleImage(Image* image, MagickBooleanType status; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); if ((checkAccelerateCondition(image) == MagickFalse) || - (checkPixelIntensity(image,method) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) + (checkPixelIntensity(image,method) == MagickFalse)) return(MagickFalse); if (image->number_channels < 3) @@ -3655,7 +3572,11 @@ MagickExport MagickBooleanType AccelerateGrayscaleImage(Image* image, (GetPixelBlueTraits(image) == UndefinedPixelTrait)) return(MagickFalse); - status=ComputeGrayscaleImage(image,method,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return(MagickFalse); + + status=ComputeGrayscaleImage(image,clEnv,method,exception); return(status); } @@ -3671,7 +3592,7 @@ MagickExport MagickBooleanType AccelerateGrayscaleImage(Image* image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static Image *ComputeLocalContrastImage(const Image *image, +static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv, const double radius,const double strength,ExceptionInfo *exception) { CacheView @@ -3681,9 +3602,6 @@ static Image *ComputeLocalContrastImage(const Image *image, cl_command_queue queue; - cl_context - context; - cl_int clStatus, iRadius; @@ -3713,8 +3631,8 @@ static Image *ComputeLocalContrastImage(const Image *image, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -3729,10 +3647,8 @@ static Image *ComputeLocalContrastImage(const Image *image, imageRows, passes; - clEnv = NULL; filteredImage = NULL; filteredImage_view = NULL; - context = NULL; imageBuffer = NULL; filteredImageBuffer = NULL; tempImageBuffer = NULL; @@ -3742,17 +3658,16 @@ static Image *ComputeLocalContrastImage(const Image *image, queue = NULL; outputReady = MagickFalse; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); /* Create and initialize OpenCL buffers. */ { - image_view=AcquireVirtualCacheView(image,exception); - inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); + image_view=AcquireAuthenticCacheView(image,exception); + inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); if (inputPixels == (const void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); goto cleanup; } @@ -3769,10 +3684,10 @@ static Image *ComputeLocalContrastImage(const Image *image, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); + imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } } @@ -3783,14 +3698,14 @@ static Image *ComputeLocalContrastImage(const Image *image, assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); if (filteredPixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -3807,10 +3722,10 @@ static Image *ComputeLocalContrastImage(const Image *image, /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); + filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } } @@ -3819,27 +3734,27 @@ static Image *ComputeLocalContrastImage(const Image *image, /* create temp buffer */ { length = image->columns * image->rows; - tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(float), NULL, &clStatus); + tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } } /* get the opencl kernel */ { - blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "LocalContrastBlurRow"); + blurRowKernel = AcquireOpenCLKernel(device,"LocalContrastBlurRow"); if (blurRowKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; - blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "LocalContrastBlurApplyColumn"); + blurColumnKernel = AcquireOpenCLKernel(device,"LocalContrastBlurApplyColumn"); if (blurColumnKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; } @@ -3863,7 +3778,7 @@ static Image *ComputeLocalContrastImage(const Image *image, if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -3886,11 +3801,10 @@ static Image *ComputeLocalContrastImage(const Image *image, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, goffset, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,LocalContrastBlurRowKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,LocalContrastBlurRowKernel,event); } } @@ -3907,7 +3821,7 @@ static Image *ComputeLocalContrastImage(const Image *image, if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -3930,11 +3844,10 @@ static Image *ComputeLocalContrastImage(const Image *image, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, goffset, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - RecordProfileData(clEnv,LocalContrastBlurApplyColumnKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,LocalContrastBlurApplyColumnKernel,event); } } } @@ -3952,14 +3865,13 @@ static Image *ComputeLocalContrastImage(const Image *image, } if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (filteredImage_view != NULL) @@ -3969,9 +3881,9 @@ cleanup: if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer); if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); - if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel); - if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (blurRowKernel!=NULL) RelinquishOpenCLKernel(blurRowKernel); + if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(blurColumnKernel); + if (queue != NULL) RelinquishOpenCLCommandQueue(device, queue); if (outputReady == MagickFalse) { if (filteredImage != NULL) @@ -3989,14 +3901,21 @@ MagickExport Image *AccelerateLocalContrastImage(const Image *image, Image *filteredImage; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateConditionRGBA(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return NULL; + if (checkAccelerateConditionRGBA(image) == MagickFalse) + return((Image *) NULL); - filteredImage=ComputeLocalContrastImage(image,radius,strength,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return((Image *) NULL); + + filteredImage=ComputeLocalContrastImage(image,clEnv,radius,strength, + exception); return(filteredImage); } @@ -4012,7 +3931,7 @@ MagickExport Image *AccelerateLocalContrastImage(const Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static MagickBooleanType ComputeModulateImage(Image *image, +static MagickBooleanType ComputeModulateImage(Image *image,MagickCLEnv clEnv, const double percent_brightness,const double percent_hue, const double percent_saturation,const ColorspaceType colorspace, ExceptionInfo *exception) @@ -4025,9 +3944,6 @@ static MagickBooleanType ComputeModulateImage(Image *image, hue, saturation; - cl_context - context; - cl_command_queue queue; @@ -4050,8 +3966,8 @@ static MagickBooleanType ComputeModulateImage(Image *image, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -4074,9 +3990,8 @@ static MagickBooleanType ComputeModulateImage(Image *image, /* * initialize opencl env */ - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); outputReady = MagickFalse; @@ -4088,7 +4003,7 @@ static MagickBooleanType ComputeModulateImage(Image *image, inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); if (inputPixels == (void *) NULL) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); goto cleanup; } @@ -4106,17 +4021,17 @@ static MagickBooleanType ComputeModulateImage(Image *image, } /* create a CL buffer from image pixel buffer */ length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); + imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - modulateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Modulate"); + modulateKernel = AcquireOpenCLKernel(device, "Modulate"); if (modulateKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -4133,7 +4048,7 @@ static MagickBooleanType ComputeModulateImage(Image *image, clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&color); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -4145,12 +4060,10 @@ static MagickBooleanType ComputeModulateImage(Image *image, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,ModulateKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,ModulateKernel,event); } if (ALIGNED(inputPixels,CLPixelPacket)) @@ -4165,23 +4078,22 @@ static MagickBooleanType ComputeModulateImage(Image *image, } if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady=SyncCacheViewAuthenticPixels(image_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (modulateKernel!=NULL) - RelinquishOpenCLKernel(clEnv, modulateKernel); + RelinquishOpenCLKernel(modulateKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(clEnv, queue); + RelinquishOpenCLCommandQueue(device,queue); return outputReady; @@ -4195,17 +4107,23 @@ MagickExport MagickBooleanType AccelerateModulateImage(Image *image, MagickBooleanType status; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateConditionRGBA(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) + if (checkAccelerateConditionRGBA(image) == MagickFalse) return(MagickFalse); if ((colorspace != HSLColorspace) && (colorspace != UndefinedColorspace)) return(MagickFalse); - status=ComputeModulateImage(image,percent_brightness,percent_hue, + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return(MagickFalse); + + status=ComputeModulateImage(image,clEnv,percent_brightness,percent_hue, percent_saturation,colorspace,exception); return(status); } @@ -4222,8 +4140,9 @@ MagickExport MagickBooleanType AccelerateModulateImage(Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, - const size_t width,const OffsetInfo *offset,ExceptionInfo *exception) +static Image* ComputeMotionBlurImage(const Image *image,MagickCLEnv clEnv, + const double *kernel,const size_t width,const OffsetInfo *offset, + ExceptionInfo *exception) { CacheView *filteredImage_view, @@ -4232,9 +4151,6 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, cl_command_queue queue; - cl_context - context; - cl_float4 biasPixel; @@ -4271,8 +4187,8 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; PixelInfo bias; @@ -4295,7 +4211,6 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, *hostPtr; outputReady = MagickFalse; - context = NULL; filteredImage = NULL; filteredImage_view = NULL; imageBuffer = NULL; @@ -4304,13 +4219,12 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, motionBlurKernel = NULL; queue = NULL; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); + device = GetOpenCLDevice(clEnv); /* Create and initialize OpenCL buffers. */ - image_view=AcquireVirtualCacheView(image,exception); - inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); + image_view=AcquireAuthenticCacheView(image,exception); + inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); if (inputPixels == (const void *) NULL) { (void) ThrowMagickException(exception,GetMagickModule(),CacheError, @@ -4331,7 +4245,7 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, } // create a CL buffer from image pixel buffer length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, + imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { @@ -4371,7 +4285,7 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, } // create a CL buffer from image pixel buffer length = image->columns * image->rows; - filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, + filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { @@ -4381,7 +4295,7 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, } - imageKernelBuffer = clEnv->library->clCreateBuffer(context, + imageKernelBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) @@ -4391,7 +4305,7 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, goto cleanup; } - queue = AcquireOpenCLCommandQueue(clEnv); + queue = AcquireOpenCLCommandQueue(device); kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, width * sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) @@ -4413,7 +4327,7 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, goto cleanup; } - offsetBuffer = clEnv->library->clCreateBuffer(context, + offsetBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(cl_int2), NULL, &clStatus); if (clStatus != CL_SUCCESS) @@ -4447,8 +4361,7 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, // get the OpenCL kernel - motionBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, - "MotionBlur"); + motionBlurKernel = AcquireOpenCLKernel(device,"MotionBlur"); if (motionBlurKernel == NULL) { (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, @@ -4508,9 +4421,7 @@ static Image* ComputeMotionBlurImage(const Image *image,const double *kernel, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,MotionBlurKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,MotionBlurKernel,event); if (ALIGNED(filteredPixels,CLPixelPacket)) { @@ -4542,8 +4453,8 @@ cleanup: if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); - if (motionBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, motionBlurKernel); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (motionBlurKernel!=NULL) RelinquishOpenCLKernel(motionBlurKernel); + if (queue != NULL) RelinquishOpenCLCommandQueue(device,queue); if (outputReady == MagickFalse && filteredImage != NULL) filteredImage=DestroyImage(filteredImage); @@ -4557,240 +4468,26 @@ MagickExport Image *AccelerateMotionBlurImage(const Image *image, Image *filteredImage; + MagickCLEnv + clEnv; + assert(image != NULL); assert(kernel != (double *) NULL); assert(offset != (OffsetInfo *) NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateConditionRGBA(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return NULL; + if (checkAccelerateConditionRGBA(image) == MagickFalse) + return((Image *) NULL); - filteredImage=ComputeMotionBlurImage(image,kernel,width,offset,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return((Image *) NULL); + + filteredImage=ComputeMotionBlurImage(image,clEnv,kernel,width,offset, + exception); return(filteredImage); } -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -% A c c e l e r a t e R a n d o m I m a g e % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -*/ - -static MagickBooleanType LaunchRandomImageKernel(MagickCLEnv clEnv, - cl_command_queue queue,cl_mem imageBuffer,const unsigned int imageColumns, - const unsigned int imageRows,cl_mem seedBuffer, - const unsigned int numGenerators,ExceptionInfo *exception) -{ - int - k; - - cl_int - clStatus; - - cl_kernel - randomImageKernel; - - cl_event - event; - - MagickBooleanType - status; - - size_t - global_work_size, - local_work_size; - - status = MagickFalse; - randomImageKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RandomNumberGenerator"); - - k = 0; - clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_mem),(void*)&imageBuffer); - clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_uint),(void*)&imageColumns); - clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_uint),(void*)&imageRows); - clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_mem),(void*)&seedBuffer); - { - const float randNormNumerator = 1.0f; - const unsigned int randNormDenominator = (unsigned int)(~0UL); - clEnv->library->clSetKernelArg(randomImageKernel,k++, - sizeof(float),(void*)&randNormNumerator); - clEnv->library->clSetKernelArg(randomImageKernel,k++, - sizeof(cl_uint),(void*)&randNormDenominator); - } - - - global_work_size = numGenerators; - local_work_size = 64; - - clStatus = clEnv->library->clEnqueueNDRangeKernel(queue,randomImageKernel,1,NULL,&global_work_size, - &local_work_size, 0, NULL, &event); - - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, - "clEnv->library->clSetKernelArg failed.", "'%s'", "."); - goto cleanup; - } - RecordProfileData(clEnv,RandomNumberGeneratorKernel,event); - clEnv->library->clReleaseEvent(event); - - status = MagickTrue; - -cleanup: - if (randomImageKernel!=NULL) RelinquishOpenCLKernel(clEnv, randomImageKernel); - return(status); -} - -static MagickBooleanType ComputeRandomImage(Image* image, - ExceptionInfo* exception) -{ - CacheView - *image_view; - - cl_command_queue - queue; - - cl_context - context; - - cl_int - clStatus; - - /* Don't release this buffer in this function !!! */ - cl_mem - randomNumberSeedsBuffer; - - cl_mem_flags - mem_flags; - - cl_mem - imageBuffer; - - MagickBooleanType - outputReady, - status; - - MagickCLEnv - clEnv; - - MagickSizeType - length; - - void - *inputPixels; - - status = MagickFalse; - outputReady = MagickFalse; - inputPixels = NULL; - context = NULL; - imageBuffer = NULL; - queue = NULL; - - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - - /* Create and initialize OpenCL buffers. */ - image_view=AcquireAuthenticCacheView(image,exception); - inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); - if (inputPixels == (void *) NULL) - { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); - goto cleanup; - } - - /* If the host pointer is aligned to the size of CLPixelPacket, - then use the host buffer directly from the GPU; otherwise, - create a buffer on the GPU and copy the data over */ - if (ALIGNED(inputPixels,CLPixelPacket)) - { - mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; - } - else - { - mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; - } - /* create a CL buffer from image pixel buffer */ - length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); - goto cleanup; - } - - queue = AcquireOpenCLCommandQueue(clEnv); - - randomNumberSeedsBuffer = GetAndLockRandSeedBuffer(clEnv); - if (randomNumberSeedsBuffer==NULL) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), - ResourceLimitWarning, "Failed to get GPU random number generators.", - "'%s'", "."); - goto cleanup; - } - - status = LaunchRandomImageKernel(clEnv,queue, - imageBuffer, - (unsigned int) image->columns, - (unsigned int) image->rows, - randomNumberSeedsBuffer, - GetNumRandGenerators(clEnv), - exception); - if (status==MagickFalse) - { - goto cleanup; - } - - if (ALIGNED(inputPixels,CLPixelPacket)) - { - length = image->columns * image->rows; - clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); - } - else - { - length = image->columns * image->rows; - clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL); - } - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); - goto cleanup; - } - outputReady=SyncCacheViewAuthenticPixels(image_view,exception); - -cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); - - image_view=DestroyCacheView(image_view); - - UnlockRandSeedBuffer(clEnv); - if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); - return outputReady; -} - -MagickExport MagickBooleanType AccelerateRandomImage(Image *image, - ExceptionInfo* exception) -{ - MagickBooleanType - status; - - assert(image != NULL); - assert(exception != (ExceptionInfo *) NULL); - - if ((checkAccelerateConditionRGBA(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return(MagickFalse); - - status=ComputeRandomImage(image,exception); - return(status); -} - /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % @@ -4804,10 +4501,11 @@ MagickExport MagickBooleanType AccelerateRandomImage(Image *image, */ static MagickBooleanType resizeHorizontalFilter(MagickCLEnv clEnv, - cl_command_queue queue,cl_mem image,cl_uint number_channels,cl_uint columns, - cl_uint rows,cl_mem resizedImage,cl_uint resizedColumns,cl_uint resizedRows, - const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients, - const float xFactor,ExceptionInfo *exception) + MagickCLDevice device,cl_command_queue queue,cl_mem image, + cl_uint number_channels,cl_uint columns,cl_uint rows,cl_mem resizedImage, + cl_uint resizedColumns,cl_uint resizedRows,const ResizeFilter *resizeFilter, + cl_mem resizeFilterCubicCoefficients,const float xFactor, + ExceptionInfo *exception) { cl_kernel horizontalKernel; @@ -4839,7 +4537,6 @@ static MagickBooleanType resizeHorizontalFilter(MagickCLEnv clEnv, status = MagickFalse; size_t - deviceLocalMemorySize, gammaAccumulatorLocalMemorySize, global_work_size[2], imageCacheLocalMemorySize, @@ -4883,9 +4580,6 @@ static MagickBooleanType resizeHorizontalFilter(MagickCLEnv clEnv, pixelPerWorkgroup = workgroupSize; } - /* get the local memory size supported by the device */ - deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv); - DisableMSCWarning(4127) while(1) RestoreMSCWarning @@ -4912,7 +4606,7 @@ RestoreMSCWarning gammaAccumulatorLocalMemorySize = sizeof(float); totalLocalMemorySize+=gammaAccumulatorLocalMemorySize; - if (totalLocalMemorySize <= deviceLocalMemorySize) + if (totalLocalMemorySize <= device->local_memory_size) break; else { @@ -4930,10 +4624,10 @@ RestoreMSCWarning resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter); resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter); - horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilter"); + horizontalKernel = AcquireOpenCLKernel(device, "ResizeHorizontalFilter"); if (horizontalKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -4976,7 +4670,7 @@ RestoreMSCWarning if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -4989,28 +4683,26 @@ RestoreMSCWarning (void) local_work_size; if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,ResizeHorizontalKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,ResizeHorizontalKernel,event); status = MagickTrue; cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); - if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel); + if (horizontalKernel != NULL) RelinquishOpenCLKernel(horizontalKernel); return(status); } static MagickBooleanType resizeVerticalFilter(MagickCLEnv clEnv, - cl_command_queue queue,cl_mem image,cl_uint number_channels,cl_uint columns, - cl_uint rows,cl_mem resizedImage,cl_uint resizedColumns,cl_uint resizedRows, - const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients, - const float yFactor,ExceptionInfo *exception) + MagickCLDevice device,cl_command_queue queue,cl_mem image, + cl_uint number_channels,cl_uint columns,cl_uint rows,cl_mem resizedImage, + cl_uint resizedColumns,cl_uint resizedRows,const ResizeFilter *resizeFilter, + cl_mem resizeFilterCubicCoefficients,const float yFactor, + ExceptionInfo *exception) { cl_kernel verticalKernel; @@ -5042,7 +4734,6 @@ static MagickBooleanType resizeVerticalFilter(MagickCLEnv clEnv, status = MagickFalse; size_t - deviceLocalMemorySize, gammaAccumulatorLocalMemorySize, global_work_size[2], imageCacheLocalMemorySize, @@ -5086,9 +4777,6 @@ static MagickBooleanType resizeVerticalFilter(MagickCLEnv clEnv, pixelPerWorkgroup = workgroupSize; } - /* get the local memory size supported by the device */ - deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv); - DisableMSCWarning(4127) while(1) RestoreMSCWarning @@ -5115,7 +4803,7 @@ RestoreMSCWarning gammaAccumulatorLocalMemorySize = sizeof(float); totalLocalMemorySize+=gammaAccumulatorLocalMemorySize; - if (totalLocalMemorySize <= deviceLocalMemorySize) + if (totalLocalMemorySize <= device->local_memory_size) break; else { @@ -5133,10 +4821,10 @@ RestoreMSCWarning resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter); resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter); - verticalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilter"); + verticalKernel = AcquireOpenCLKernel(device,"ResizeVerticalFilter"); if (verticalKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -5179,7 +4867,7 @@ RestoreMSCWarning if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -5191,24 +4879,21 @@ RestoreMSCWarning clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, verticalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,ResizeVerticalKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,ResizeVerticalKernel,event); status = MagickTrue; cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); - if (verticalKernel != NULL) RelinquishOpenCLKernel(clEnv, verticalKernel); + if (verticalKernel != NULL) RelinquishOpenCLKernel(verticalKernel); return(status); } -static Image *ComputeResizeImage(const Image* image, +static Image *ComputeResizeImage(const Image* image,MagickCLEnv clEnv, const size_t resizedColumns,const size_t resizedRows, const ResizeFilter *resizeFilter,ExceptionInfo *exception) { @@ -5222,9 +4907,6 @@ static Image *ComputeResizeImage(const Image* image, cl_int clStatus; - cl_context - context; - cl_mem cubicCoefficientsBuffer, filteredImageBuffer, @@ -5246,8 +4928,8 @@ static Image *ComputeResizeImage(const Image* image, outputReady, status; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -5264,8 +4946,6 @@ static Image *ComputeResizeImage(const Image* image, outputReady = MagickFalse; filteredImage = NULL; filteredImage_view = NULL; - clEnv = NULL; - context = NULL; imageBuffer = NULL; tempImageBuffer = NULL; filteredImageBuffer = NULL; @@ -5273,26 +4953,25 @@ static Image *ComputeResizeImage(const Image* image, cubicCoefficientsBuffer = NULL; queue = NULL; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); + device = GetOpenCLDevice(clEnv); - image_view = AcquireVirtualCacheView(image, exception); - imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); + image_view = AcquireAuthenticCacheView(image, exception); + imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; - cubicCoefficientsBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, 7 * sizeof(float), NULL, &clStatus); + cubicCoefficientsBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY, 7 * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - queue = AcquireOpenCLCommandQueue(clEnv); + queue = AcquireOpenCLCommandQueue(device); mappedCoefficientBuffer = (float*)clEnv->library->clEnqueueMapBuffer(queue, cubicCoefficientsBuffer, CL_TRUE, CL_MAP_WRITE, 0, 7 * sizeof(float) , 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); goto cleanup; } resizeFilterCoefficient = GetResizeFilterCoefficient(resizeFilter); @@ -5303,7 +4982,7 @@ static Image *ComputeResizeImage(const Image* image, clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, cubicCoefficientsBuffer, mappedCoefficientBuffer, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } @@ -5312,12 +4991,12 @@ static Image *ComputeResizeImage(const Image* image, goto cleanup; if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception); filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, - context,filteredPixels,exception); + device,filteredPixels,exception); if (filteredImageBuffer == (cl_mem) NULL) goto cleanup; @@ -5327,21 +5006,21 @@ static Image *ComputeResizeImage(const Image* image, if (xFactor > yFactor) { length = resizedColumns*image->rows*number_channels; - tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLQuantum), NULL, &clStatus); + tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length*sizeof(CLQuantum), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - status = resizeHorizontalFilter(clEnv,queue,imageBuffer,number_channels, + status = resizeHorizontalFilter(clEnv,device,queue,imageBuffer,number_channels, (cl_uint) image->columns,(cl_uint) image->rows,tempImageBuffer, (cl_uint) resizedColumns,(cl_uint) image->rows,resizeFilter, cubicCoefficientsBuffer,xFactor,exception); if (status != MagickTrue) goto cleanup; - status = resizeVerticalFilter(clEnv,queue,tempImageBuffer,number_channels, + status = resizeVerticalFilter(clEnv,device,queue,tempImageBuffer,number_channels, (cl_uint) resizedColumns,(cl_uint) image->rows,filteredImageBuffer, (cl_uint) resizedColumns,(cl_uint) resizedRows,resizeFilter, cubicCoefficientsBuffer,yFactor,exception); @@ -5351,21 +5030,21 @@ static Image *ComputeResizeImage(const Image* image, else { length = image->columns*resizedRows*number_channels; - tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLQuantum), NULL, &clStatus); + tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length*sizeof(CLQuantum), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - status = resizeVerticalFilter(clEnv,queue,imageBuffer,number_channels, + status = resizeVerticalFilter(clEnv,device,queue,imageBuffer,number_channels, (cl_uint) image->columns,(cl_int) image->rows,tempImageBuffer, (cl_uint) image->columns,(cl_uint) resizedRows,resizeFilter, cubicCoefficientsBuffer,yFactor,exception); if (status != MagickTrue) goto cleanup; - status = resizeHorizontalFilter(clEnv,queue,tempImageBuffer,number_channels, + status = resizeHorizontalFilter(clEnv,device,queue,tempImageBuffer,number_channels, (cl_uint) image->columns, (cl_uint) resizedRows,filteredImageBuffer, (cl_uint) resizedColumns, (cl_uint) resizedRows,resizeFilter, cubicCoefficientsBuffer,xFactor,exception); @@ -5373,13 +5052,12 @@ static Image *ComputeResizeImage(const Image* image, goto cleanup; } - if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) + if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) goto cleanup; outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (filteredImage_view != NULL) @@ -5389,7 +5067,7 @@ cleanup: if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer); if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); if (cubicCoefficientsBuffer!=NULL) clEnv->library->clReleaseMemObject(cubicCoefficientsBuffer); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (queue != NULL) RelinquishOpenCLCommandQueue(device, queue); if (outputReady == MagickFalse && filteredImage != NULL) filteredImage=DestroyImage(filteredImage); return(filteredImage); @@ -5418,20 +5096,26 @@ MagickExport Image *AccelerateResizeImage(const Image *image, Image *filteredImage; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateCondition(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return NULL; + if (checkAccelerateCondition(image) == MagickFalse) + return((Image *) NULL); if ((gpuSupportedResizeWeighting(GetResizeFilterWeightingType( resizeFilter)) == MagickFalse) || (gpuSupportedResizeWeighting(GetResizeFilterWindowWeightingType( resizeFilter)) == MagickFalse)) - return NULL; + return((Image *) NULL); - filteredImage=ComputeResizeImage(image,resizedColumns,resizedRows, + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return((Image *) NULL); + + filteredImage=ComputeResizeImage(image,clEnv,resizedColumns,resizedRows, resizeFilter,exception); return(filteredImage); } @@ -5448,8 +5132,8 @@ MagickExport Image *AccelerateResizeImage(const Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static Image* ComputeRotationalBlurImage(const Image *image,const double angle, - ExceptionInfo *exception) +static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv, + const double angle,ExceptionInfo *exception) { CacheView *image_view, @@ -5458,9 +5142,6 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, cl_command_queue queue; - cl_context - context; - cl_float2 blurCenter; @@ -5499,8 +5180,8 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; PixelInfo bias; @@ -5515,7 +5196,6 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, *filteredPixels; outputReady = MagickFalse; - context = NULL; filteredImage = NULL; filteredImage_view = NULL; filteredPixels = NULL; @@ -5526,11 +5206,10 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, queue = NULL; rotationalBlurKernel = NULL; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); + device = GetOpenCLDevice(clEnv); - image_view=AcquireVirtualCacheView(image, exception); - imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); + image_view=AcquireAuthenticCacheView(image, exception); + imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -5539,12 +5218,12 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, goto cleanup; if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue) { - (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception); filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, - context,filteredPixels,exception); + device,filteredPixels,exception); if (filteredImageBuffer == (cl_mem) NULL) goto cleanup; @@ -5554,32 +5233,32 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, cossin_theta_size=(unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt((double)blurRadius)+2UL); /* create a buffer for sin_theta and cos_theta */ - sinThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus); + sinThetaBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - cosThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus); + cosThetaBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } - queue = AcquireOpenCLCommandQueue(clEnv); + queue = AcquireOpenCLCommandQueue(device); sinThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, sinThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); goto cleanup; } cosThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, cosThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); goto cleanup; } @@ -5595,15 +5274,15 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, clStatus |= clEnv->library->clEnqueueUnmapMemObject(queue, cosThetaBuffer, cosThetaPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } /* get the OpenCL kernel */ - rotationalBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RotationalBlur"); + rotationalBlurKernel = AcquireOpenCLKernel(device,"RotationalBlur"); if (rotationalBlurKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -5628,7 +5307,7 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -5638,23 +5317,20 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, rotationalBlurKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,RotationalBlurKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,RotationalBlurKernel,event); - if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) + if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) { - (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (filteredImage_view != NULL) @@ -5664,8 +5340,8 @@ cleanup: if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (sinThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(sinThetaBuffer); if (cosThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(cosThetaBuffer); - if (rotationalBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, rotationalBlurKernel); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (rotationalBlurKernel!=NULL) RelinquishOpenCLKernel(rotationalBlurKernel); + if (queue != NULL) RelinquishOpenCLCommandQueue(device,queue); if (outputReady == MagickFalse) { if (filteredImage != NULL) @@ -5683,14 +5359,20 @@ MagickExport Image* AccelerateRotationalBlurImage(const Image *image, Image *filteredImage; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateCondition(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return NULL; + if (checkAccelerateCondition(image) == MagickFalse) + return((Image *) NULL); - filteredImage=ComputeRotationalBlurImage(image,angle,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return((Image *) NULL); + + filteredImage=ComputeRotationalBlurImage(image,clEnv,angle,exception); return filteredImage; } @@ -5706,9 +5388,9 @@ MagickExport Image* AccelerateRotationalBlurImage(const Image *image, %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% */ -static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, - const double sigma,const double gain,const double threshold, - ExceptionInfo *exception) +static Image *ComputeUnsharpMaskImage(const Image *image,MagickCLEnv clEnv, + const double radius,const double sigma,const double gain, + const double threshold,ExceptionInfo *exception) { CacheView *filteredImage_view, @@ -5717,9 +5399,6 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -5755,8 +5434,8 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; MagickSizeType length; @@ -5767,10 +5446,8 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, unsigned int i; - clEnv = NULL; filteredImage = NULL; filteredImage_view = NULL; - context = NULL; imageBuffer = NULL; filteredImageBuffer = NULL; filteredPixels = NULL; @@ -5781,12 +5458,11 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, queue = NULL; outputReady = MagickFalse; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); - image_view = AcquireVirtualCacheView(image, exception); - imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); + image_view = AcquireAuthenticCacheView(image, exception); + imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -5795,43 +5471,43 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, goto cleanup; if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue) { - (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception); filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, - context,filteredPixels,exception); + device,filteredPixels,exception); if (filteredImageBuffer == (cl_mem) NULL) goto cleanup; - imageKernelBuffer=createKernelInfo(clEnv,context,queue,radius,sigma, + imageKernelBuffer=createKernelInfo(clEnv,device,queue,radius,sigma, &kernelWidth,exception); { /* create temp buffer */ { length = image->columns * image->rows; - tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(cl_float4), NULL, &clStatus); + tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length * sizeof(cl_float4), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); goto cleanup; } } /* get the opencl kernel */ { - blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow"); + blurRowKernel = AcquireOpenCLKernel(device,"BlurRow"); if (blurRowKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; - unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumn"); + unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(device,"UnsharpMaskBlurColumn"); if (unsharpMaskBlurColumnKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; } @@ -5856,7 +5532,7 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -5874,12 +5550,10 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,BlurRowKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,BlurRowKernel,event); } @@ -5905,7 +5579,7 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -5923,27 +5597,24 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,UnsharpMaskBlurColumnKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,UnsharpMaskBlurColumnKernel,event); } } /* get result */ - if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) + if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) { - (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (filteredImage_view != NULL) @@ -5953,9 +5624,9 @@ cleanup: if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer); if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); - if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel); - if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (blurRowKernel!=NULL) RelinquishOpenCLKernel(blurRowKernel); + if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(unsharpMaskBlurColumnKernel); + if (queue != NULL) RelinquishOpenCLCommandQueue(device,queue); if (outputReady == MagickFalse) { if (filteredImage != NULL) @@ -5968,8 +5639,8 @@ cleanup: } static Image *ComputeUnsharpMaskImageSingle(const Image *image, - const double radius,const double sigma,const double gain, - const double threshold,int blurOnly,ExceptionInfo *exception) + MagickCLEnv clEnv,const double radius,const double sigma,const double gain, + const double threshold,ExceptionInfo *exception) { CacheView *filteredImage_view, @@ -5978,11 +5649,7 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, cl_command_queue queue; - cl_context - context; - cl_int - justBlur, clStatus; cl_kernel @@ -6013,17 +5680,15 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; void *filteredPixels; - clEnv = NULL; filteredImage = NULL; filteredImage_view = NULL; filteredPixels = NULL; - context = NULL; imageBuffer = NULL; filteredImageBuffer = NULL; imageKernelBuffer = NULL; @@ -6031,12 +5696,11 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, queue = NULL; outputReady = MagickFalse; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); - image_view=AcquireVirtualCacheView(image,exception); - imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); + image_view=AcquireAuthenticCacheView(image,exception); + imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -6046,20 +5710,20 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, - context,filteredPixels,exception); + device,filteredPixels,exception); if (filteredImageBuffer == (void *) NULL) goto cleanup; - imageKernelBuffer=createKernelInfo(clEnv,context,queue,radius,sigma, + imageKernelBuffer=createKernelInfo(clEnv,device,queue,radius,sigma, &kernelWidth,exception); { /* get the opencl kernel */ { - unsharpMaskKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMask"); + unsharpMaskKernel = AcquireOpenCLKernel(device, "UnsharpMask"); if (unsharpMaskKernel == NULL) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; } @@ -6070,7 +5734,6 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, number_channels = (cl_uint) image->number_channels; fGain = (float) gain; fThreshold = (float) threshold; - justBlur = blurOnly; /* set the kernel arguments */ i = 0; @@ -6084,11 +5747,10 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_float4)*(8 * (32 + kernelWidth)),(void *) NULL); clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fGain); clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fThreshold); - clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&justBlur); clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -6106,22 +5768,19 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, clStatus = clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskKernel, 2, NULL, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } - clEnv->library->clFlush(queue); - RecordProfileData(clEnv,UnsharpMaskKernel,event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,UnsharpMaskKernel,event); } } - if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) + if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) goto cleanup; outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); image_view=DestroyCacheView(image_view); if (filteredImage_view != NULL) @@ -6130,8 +5789,8 @@ cleanup: if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer); - if (unsharpMaskKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskKernel); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (unsharpMaskKernel!=NULL) RelinquishOpenCLKernel(unsharpMaskKernel); + if (queue != NULL) RelinquishOpenCLCommandQueue(device,queue); if (outputReady == MagickFalse) { if (filteredImage != NULL) @@ -6150,23 +5809,29 @@ MagickExport Image *AccelerateUnsharpMaskImage(const Image *image, Image *filteredImage; + MagickCLEnv + clEnv; + assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateCondition(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return NULL; + if (checkAccelerateCondition(image) == MagickFalse) + return((Image *) NULL); + + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return((Image *) NULL); if (radius < 12.1) - filteredImage=ComputeUnsharpMaskImageSingle(image,radius,sigma,gain, - threshold,0,exception); + filteredImage=ComputeUnsharpMaskImageSingle(image,clEnv,radius,sigma,gain, + threshold,exception); else - filteredImage=ComputeUnsharpMaskImage(image,radius,sigma,gain,threshold, - exception); + filteredImage=ComputeUnsharpMaskImage(image,clEnv,radius,sigma,gain, + threshold,exception); return(filteredImage); } -static Image *ComputeWaveletDenoiseImage(const Image *image, +static Image *ComputeWaveletDenoiseImage(const Image *image,MagickCLEnv clEnv, const double threshold,ExceptionInfo *exception) { CacheView @@ -6176,9 +5841,6 @@ static Image *ComputeWaveletDenoiseImage(const Image *image, cl_command_queue queue; - cl_context - context; - cl_int clStatus; @@ -6198,8 +5860,8 @@ static Image *ComputeWaveletDenoiseImage(const Image *image, MagickBooleanType outputReady; - MagickCLEnv - clEnv; + MagickCLDevice + device; void *filteredPixels; @@ -6214,13 +5876,12 @@ static Image *ComputeWaveletDenoiseImage(const Image *image, denoiseKernel = NULL; outputReady = MagickFalse; - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); + device = GetOpenCLDevice(clEnv); + queue = AcquireOpenCLCommandQueue(device); /* Create and initialize OpenCL buffers. */ - image_view = AcquireVirtualCacheView(image, exception); - imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); + image_view = AcquireAuthenticCacheView(image, exception); + imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -6230,20 +5891,20 @@ static Image *ComputeWaveletDenoiseImage(const Image *image, goto cleanup; if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue) { - (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception); filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, - context,filteredPixels,exception); + device,filteredPixels,exception); if (filteredImageBuffer == (cl_mem) NULL) goto cleanup; /* get the opencl kernel */ - denoiseKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "WaveletDenoise"); + denoiseKernel = AcquireOpenCLKernel(device,"WaveletDenoise"); if (denoiseKernel == NULL) { - (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -6270,7 +5931,7 @@ static Image *ComputeWaveletDenoiseImage(const Image *image, clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_uint), (void *)&height); if (clStatus != CL_SUCCESS) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -6290,24 +5951,22 @@ static Image *ComputeWaveletDenoiseImage(const Image *image, clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, denoiseKernel, 2, NULL, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { - (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } } - RecordProfileData(clEnv, WaveletDenoiseKernel, event); - clEnv->library->clReleaseEvent(event); + RecordProfileData(device,WaveletDenoiseKernel,event); } - if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) + if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) { - (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = SyncCacheViewAuthenticPixels(filteredImage_view, exception); cleanup: - OpenCLLogException(__FUNCTION__, __LINE__, exception); image_view = DestroyCacheView(image_view); if (filteredImage_view != NULL) @@ -6315,8 +5974,8 @@ cleanup: if (imageBuffer != NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (filteredImageBuffer != NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); - if (denoiseKernel != NULL) RelinquishOpenCLKernel(clEnv, denoiseKernel); - if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); + if (denoiseKernel != NULL) RelinquishOpenCLKernel(denoiseKernel); + if (queue != NULL) RelinquishOpenCLCommandQueue(device,queue); if (outputReady == MagickFalse) { if (filteredImage != NULL) @@ -6332,16 +5991,22 @@ MagickExport Image *AccelerateWaveletDenoiseImage(const Image *image, const double threshold,ExceptionInfo *exception) { Image - *filteredImage; + *filteredImage; + + MagickCLEnv + clEnv; assert(image != NULL); assert(exception != (ExceptionInfo *)NULL); - if ((checkAccelerateCondition(image) == MagickFalse) || - (checkOpenCLEnvironment(exception) == MagickFalse)) - return (Image *) NULL; + if (checkAccelerateCondition(image) == MagickFalse) + return((Image *) NULL); - filteredImage=ComputeWaveletDenoiseImage(image,threshold,exception); + clEnv=getOpenCLEnvironment(exception); + if (clEnv == (MagickCLEnv) NULL) + return((Image *) NULL); + + filteredImage=ComputeWaveletDenoiseImage(image,clEnv,threshold,exception); return(filteredImage); } @@ -6516,7 +6181,7 @@ MagickExport MagickBooleanType AccelerateRandomImage( magick_unreferenced(image); magick_unreferenced(exception); - return MagickFalse; + return(MagickFalse); } MagickExport Image *AccelerateResizeImage(const Image *magick_unused(image), diff --git a/MagickCore/accelerate.h b/MagickCore/accelerate.h index 19e5c7db52..cfd2c5bf5d 100644 --- a/MagickCore/accelerate.h +++ b/MagickCore/accelerate.h @@ -56,8 +56,7 @@ extern MagickExport MagickBooleanType AccelerateGrayscaleImage(Image *,const PixelIntensityMethod, ExceptionInfo *), AccelerateModulateImage(Image *,const double,const double,const double, - const ColorspaceType, ExceptionInfo*), - AccelerateRandomImage(Image*,ExceptionInfo*); + const ColorspaceType, ExceptionInfo*); #if defined(__cplusplus) || defined(c_plusplus) } diff --git a/MagickCore/magick.c b/MagickCore/magick.c index 5625e288cc..8e1959d71c 100644 --- a/MagickCore/magick.c +++ b/MagickCore/magick.c @@ -1528,7 +1528,9 @@ MagickExport void MagickCoreTerminus(void) AnnotateComponentTerminus(); MimeComponentTerminus(); TypeComponentTerminus(); +#if defined(MAGICKCORE_OPENCL_SUPPORT) OpenCLTerminus(); +#endif ColorComponentTerminus(); #if defined(MAGICKCORE_WINDOWS_SUPPORT) NTWindowsTerminus(); diff --git a/MagickCore/opencl-private.h b/MagickCore/opencl-private.h index 6558728bd4..58ec7c767d 100644 --- a/MagickCore/opencl-private.h +++ b/MagickCore/opencl-private.h @@ -23,11 +23,41 @@ Include declarations. */ #include "MagickCore/studio.h" #include "MagickCore/opencl.h" +#include "MagickCore/thread_.h" #if defined(__cplusplus) || defined(c_plusplus) extern "C" { #endif +typedef enum +{ + AddNoiseKernel, + BlurColumnKernel, + BlurRowKernel, + CompositeKernel, + ContrastKernel, + ContrastStretchKernel, + ConvolveKernel, + ConvolveOptimizedKernel, + ComputeFunctionKernel, + EqualizeKernel, + GrayScaleKernel, + HistogramKernel, + HullPass1Kernel, + HullPass2Kernel, + LocalContrastBlurApplyColumnKernel, + LocalContrastBlurRowKernel, + ModulateKernel, + MotionBlurKernel, + ResizeHorizontalKernel, + ResizeVerticalKernel, + RotationalBlurKernel, + UnsharpMaskKernel, + UnsharpMaskBlurColumnKernel, + WaveletDenoiseKernel, + KERNEL_COUNT +} ProfiledKernels; + #if !defined(MAGICKCORE_OPENCL_SUPPORT) typedef void* cl_context; typedef void* cl_command_queue; @@ -36,279 +66,294 @@ extern "C" { typedef void* cl_kernel; typedef void* cl_mem; typedef void* cl_platform_id; - typedef struct { unsigned char t[8]; } cl_device_type; /* 64-bit */ + typedef void* cl_device_type; #else -#define MAX_COMMAND_QUEUES 16 - /* - * - * function pointer typedefs - * - */ + Define declarations. +*/ +#define MAGICKCORE_OPENCL_UNDEFINED_SCORE -1.0 +#define MAGICKCORE_OPENCL_PROFILE_KERNELS 0 +#define MAGICKCORE_OPENCL_COMMAND_QUEUES 16 + +#if MAGICKCORE_OPENCL_PROFILE_KERNELS +typedef struct +{ + cl_ulong min; + cl_ulong max; + cl_ulong total; + cl_ulong count; +} KernelProfileRecord; +#endif /* Platform APIs */ -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetPlatformIDs)( - cl_uint num_entries, - cl_platform_id * platforms, - cl_uint * num_platforms) CL_API_SUFFIX__VERSION_1_0; +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clGetPlatformIDs)(cl_uint num_entries, + cl_platform_id *platforms,cl_uint *num_platforms) CL_API_SUFFIX__VERSION_1_0; -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetPlatformInfo)( - cl_platform_id platform, - cl_platform_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0; - -/* Device APIs */ -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetDeviceIDs)( - cl_platform_id platform, - cl_device_type device_type, - cl_uint num_entries, - cl_device_id * devices, - cl_uint * num_devices) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetDeviceInfo)( - cl_device_id device, - cl_device_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0; - -/* Context APIs */ -typedef CL_API_ENTRY cl_context (CL_API_CALL *MAGICKpfn_clCreateContext)( - const cl_context_properties * properties, - cl_uint num_devices, - const cl_device_id * devices, - void (CL_CALLBACK *pfn_notify)(const char *, const void *, size_t, void *), - void * user_data, - cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseContext)( - cl_context context) CL_API_SUFFIX__VERSION_1_0; - -/* Command Queue APIs */ -typedef CL_API_ENTRY cl_command_queue (CL_API_CALL *MAGICKpfn_clCreateCommandQueue)( - cl_context context, - cl_device_id device, - cl_command_queue_properties properties, - cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseCommandQueue)( - cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0; - -/* Memory Object APIs */ -typedef CL_API_ENTRY cl_mem (CL_API_CALL *MAGICKpfn_clCreateBuffer)( - cl_context context, - cl_mem_flags flags, - size_t size, - void * host_ptr, - cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseMemObject)(cl_mem memobj) CL_API_SUFFIX__VERSION_1_0; - -/* Program Object APIs */ -typedef CL_API_ENTRY cl_program (CL_API_CALL *MAGICKpfn_clCreateProgramWithSource)( - cl_context context, - cl_uint count, - const char ** strings, - const size_t * lengths, - cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_program (CL_API_CALL *MAGICKpfn_clCreateProgramWithBinary)( - cl_context context, - cl_uint num_devices, - const cl_device_id * device_list, - const size_t * lengths, - const unsigned char ** binaries, - cl_int * binary_status, - cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseProgram)(cl_program program) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clBuildProgram)( - cl_program program, - cl_uint num_devices, - const cl_device_id * device_list, - const char * options, - void (CL_CALLBACK *pfn_notify)(cl_program program, void * user_data), - void * user_data) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetProgramInfo)( - cl_program program, - cl_program_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetProgramBuildInfo)( - cl_program program, - cl_device_id device, - cl_program_build_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0; - -/* Kernel Object APIs */ -typedef CL_API_ENTRY cl_kernel (CL_API_CALL *MAGICKpfn_clCreateKernel)( - cl_program program, - const char * kernel_name, - cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseKernel)(cl_kernel kernel) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clSetKernelArg)( - cl_kernel kernel, - cl_uint arg_index, - size_t arg_size, - const void * arg_value) CL_API_SUFFIX__VERSION_1_0; - -/* Flush and Finish APIs */ -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clFlush)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clFinish)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0; - -/* Enqueued Commands APIs */ -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueReadBuffer)( - cl_command_queue command_queue, - cl_mem buffer, - cl_bool blocking_read, - size_t offset, - size_t cb, - void * ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueWriteBuffer)( - cl_command_queue command_queue, - cl_mem buffer, - cl_bool blocking_write, - size_t offset, - size_t cb, - const void * ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY void * (CL_API_CALL *MAGICKpfn_clEnqueueMapBuffer)( - cl_command_queue command_queue, - cl_mem buffer, - cl_bool blocking_map, - cl_map_flags map_flags, - size_t offset, - size_t cb, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event, - cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueUnmapMemObject)( - cl_command_queue command_queue, - cl_mem memobj, - void * mapped_ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueNDRangeKernel)( - cl_command_queue command_queue, - cl_kernel kernel, - cl_uint work_dim, - const size_t * global_work_offset, - const size_t * global_work_size, - const size_t * local_work_size, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event) CL_API_SUFFIX__VERSION_1_0; - -typedef CL_API_ENTRY cl_int(CL_API_CALL *MAGICKpfn_clGetEventProfilingInfo)( - cl_event event, - cl_profiling_info param_name, - size_t param_value_size, - void *param_value, +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clGetPlatformInfo)(cl_platform_id platform, + cl_platform_info param_name,size_t param_value_size,void *param_value, size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0; -typedef CL_API_ENTRY cl_int(CL_API_CALL *MAGICKpfn_clWaitForEvents)( - cl_uint num_events, + +/* Device APIs */ +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clGetDeviceIDs)(cl_platform_id platform, + cl_device_type device_type,cl_uint num_entries,cl_device_id *devices, + cl_uint *num_devices) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clGetDeviceInfo)(cl_device_id device, + cl_device_info param_name,size_t param_value_size,void *param_value, + size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0; + + +/* Context APIs */ +typedef CL_API_ENTRY cl_context + (CL_API_CALL *MAGICKpfn_clCreateContext)( + const cl_context_properties *properties,cl_uint num_devices, + const cl_device_id *devices,void (CL_CALLBACK *pfn_notify)(const char *, + const void *,size_t,void *),void *user_data,cl_int *errcode_ret) + CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clReleaseContext)(cl_context context) + CL_API_SUFFIX__VERSION_1_0; + + +/* Command Queue APIs */ +typedef CL_API_ENTRY cl_command_queue + (CL_API_CALL *MAGICKpfn_clCreateCommandQueue)(cl_context context, + cl_device_id device,cl_command_queue_properties properties, + cl_int *errcode_ret) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clReleaseCommandQueue)( + cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0; + + +/* Memory Object APIs */ +typedef CL_API_ENTRY cl_mem + (CL_API_CALL *MAGICKpfn_clCreateBuffer)(cl_context context, + cl_mem_flags flags,size_t size,void *host_ptr,cl_int *errcode_ret) + CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clReleaseMemObject)(cl_mem memobj) + CL_API_SUFFIX__VERSION_1_0; + + +/* Program Object APIs */ +typedef CL_API_ENTRY cl_program + (CL_API_CALL *MAGICKpfn_clCreateProgramWithSource)(cl_context context, + cl_uint count,const char **strings,const size_t *lengths, + cl_int *errcode_ret) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_program + (CL_API_CALL *MAGICKpfn_clCreateProgramWithBinary)(cl_context context, + cl_uint num_devices,const cl_device_id *device_list,const size_t *lengths, + const unsigned char **binaries,cl_int *binary_status,cl_int *errcode_ret) + CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clReleaseProgram)(cl_program program) + CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clBuildProgram)(cl_program program, + cl_uint num_devices,const cl_device_id *device_list,const char *options, + void (CL_CALLBACK *pfn_notify)(cl_program program,void * user_data), + void *user_data) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clGetProgramBuildInfo)(cl_program program, + cl_device_id device,cl_program_build_info param_name,size_t param_value_size, + void *param_value,size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clGetProgramInfo)(cl_program program, + cl_program_info param_name,size_t param_value_size,void *param_value, + size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0; + + +/* Kernel Object APIs */ +typedef CL_API_ENTRY cl_kernel + (CL_API_CALL *MAGICKpfn_clCreateKernel)(cl_program program, + const char *kernel_name,cl_int *errcode_ret) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clReleaseKernel)(cl_kernel kernel) + CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clSetKernelArg)(cl_kernel kernel,cl_uint arg_index, + size_t arg_size,const void * arg_value) CL_API_SUFFIX__VERSION_1_0; + + +/* Enqueued Commands APIs */ +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clEnqueueReadBuffer)(cl_command_queue command_queue, + cl_mem buffer,cl_bool blocking_read,size_t offset,size_t cb,void *ptr, + cl_uint num_events_in_wait_list,const cl_event *event_wait_list, + cl_event *event) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY void + *(CL_API_CALL *MAGICKpfn_clEnqueueMapBuffer)(cl_command_queue command_queue, + cl_mem buffer,cl_bool blocking_map,cl_map_flags map_flags,size_t offset, + size_t cb,cl_uint num_events_in_wait_list,const cl_event *event_wait_list, + cl_event *event,cl_int *errcode_ret) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clEnqueueUnmapMemObject)( + cl_command_queue command_queue,cl_mem memobj,void *mapped_ptr, + cl_uint num_events_in_wait_list,const cl_event *event_wait_list, + cl_event *event) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clEnqueueNDRangeKernel)( + cl_command_queue command_queue,cl_kernel kernel,cl_uint work_dim, + const size_t *global_work_offset,const size_t *global_work_size, + const size_t *local_work_size,cl_uint num_events_in_wait_list, + const cl_event * event_wait_list,cl_event *event) + CL_API_SUFFIX__VERSION_1_0; + + +/* Profiling APIs */ +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clGetEventProfilingInfo)(cl_event event, + cl_profiling_info param_name,size_t param_value_size,void *param_value, + size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clWaitForEvents)(cl_uint num_events, const cl_event *event_list) CL_API_SUFFIX__VERSION_1_0; -typedef CL_API_ENTRY cl_int(CL_API_CALL *MAGICKpfn_clReleaseEvent)( - cl_event event) CL_API_SUFFIX__VERSION_1_0; +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clReleaseEvent)(cl_event event) + CL_API_SUFFIX__VERSION_1_0; -/* - * - * vendor dispatch table structure - * - * note that the types in the structure KHRicdVendorDispatch mirror the function - * names listed in the string table khrIcdVendorDispatchFunctionNames - * - */ + +/* Finish APIs, only here for GetAndLockRandSeedBuffer */ +typedef CL_API_ENTRY cl_int + (CL_API_CALL *MAGICKpfn_clFinish)(cl_command_queue command_queue) + CL_API_SUFFIX__VERSION_1_0; typedef struct MagickLibraryRec MagickLibrary; struct MagickLibraryRec { - void * base; + MAGICKpfn_clGetPlatformIDs clGetPlatformIDs; + MAGICKpfn_clGetPlatformInfo clGetPlatformInfo; - MAGICKpfn_clGetPlatformIDs clGetPlatformIDs; - MAGICKpfn_clGetPlatformInfo clGetPlatformInfo; - MAGICKpfn_clGetDeviceIDs clGetDeviceIDs; - MAGICKpfn_clGetDeviceInfo clGetDeviceInfo; - MAGICKpfn_clCreateContext clCreateContext; - MAGICKpfn_clCreateCommandQueue clCreateCommandQueue; - MAGICKpfn_clReleaseCommandQueue clReleaseCommandQueue; - MAGICKpfn_clCreateBuffer clCreateBuffer; - MAGICKpfn_clReleaseMemObject clReleaseMemObject; - MAGICKpfn_clCreateProgramWithSource clCreateProgramWithSource; - MAGICKpfn_clCreateProgramWithBinary clCreateProgramWithBinary; - MAGICKpfn_clReleaseProgram clReleaseProgram; - MAGICKpfn_clBuildProgram clBuildProgram; - MAGICKpfn_clGetProgramInfo clGetProgramInfo; - MAGICKpfn_clGetProgramBuildInfo clGetProgramBuildInfo; - MAGICKpfn_clCreateKernel clCreateKernel; - MAGICKpfn_clReleaseKernel clReleaseKernel; - MAGICKpfn_clSetKernelArg clSetKernelArg; - MAGICKpfn_clFlush clFlush; - MAGICKpfn_clFinish clFinish; - MAGICKpfn_clEnqueueReadBuffer clEnqueueReadBuffer; - MAGICKpfn_clEnqueueWriteBuffer clEnqueueWriteBuffer; - MAGICKpfn_clEnqueueMapBuffer clEnqueueMapBuffer; - MAGICKpfn_clEnqueueUnmapMemObject clEnqueueUnmapMemObject; - MAGICKpfn_clEnqueueNDRangeKernel clEnqueueNDRangeKernel; - MAGICKpfn_clGetEventProfilingInfo clGetEventProfilingInfo; - MAGICKpfn_clWaitForEvents clWaitForEvents; - MAGICKpfn_clReleaseEvent clReleaseEvent; + MAGICKpfn_clGetDeviceIDs clGetDeviceIDs; + MAGICKpfn_clGetDeviceInfo clGetDeviceInfo; + + MAGICKpfn_clCreateContext clCreateContext; + MAGICKpfn_clReleaseContext clReleaseContext; + + MAGICKpfn_clCreateCommandQueue clCreateCommandQueue; + MAGICKpfn_clReleaseCommandQueue clReleaseCommandQueue; + + MAGICKpfn_clCreateBuffer clCreateBuffer; + MAGICKpfn_clReleaseMemObject clReleaseMemObject; + + MAGICKpfn_clCreateProgramWithSource clCreateProgramWithSource; + MAGICKpfn_clCreateProgramWithBinary clCreateProgramWithBinary; + MAGICKpfn_clReleaseProgram clReleaseProgram; + MAGICKpfn_clBuildProgram clBuildProgram; + MAGICKpfn_clGetProgramBuildInfo clGetProgramBuildInfo; + MAGICKpfn_clGetProgramInfo clGetProgramInfo; + + MAGICKpfn_clCreateKernel clCreateKernel; + MAGICKpfn_clReleaseKernel clReleaseKernel; + MAGICKpfn_clSetKernelArg clSetKernelArg; + + MAGICKpfn_clEnqueueReadBuffer clEnqueueReadBuffer; + MAGICKpfn_clEnqueueMapBuffer clEnqueueMapBuffer; + MAGICKpfn_clEnqueueUnmapMemObject clEnqueueUnmapMemObject; + MAGICKpfn_clEnqueueNDRangeKernel clEnqueueNDRangeKernel; + + MAGICKpfn_clGetEventProfilingInfo clGetEventProfilingInfo; + MAGICKpfn_clWaitForEvents clWaitForEvents; + MAGICKpfn_clReleaseEvent clReleaseEvent; + + MAGICKpfn_clFinish clFinish; }; -struct _MagickCLEnv { - MagickBooleanType OpenCLInitialized; /* whether OpenCL environment is initialized. */ - MagickBooleanType OpenCLDisabled; /* whether if OpenCL has been explicitely disabled. */ +struct _MagickCLDevice +{ + char + *name, + *version; - MagickLibrary * library; + cl_command_queue + command_queues[MAGICKCORE_OPENCL_COMMAND_QUEUES]; - /*OpenCL objects */ - cl_platform_id platform; - cl_device_type deviceType; - cl_device_id device; - cl_context context; + cl_context + context; - MagickBooleanType disableProgramCache; /* disable the OpenCL program cache */ - cl_program programs[MAGICK_OPENCL_NUM_PROGRAMS]; /* one program object maps one kernel source file */ + cl_device_id + deviceID; - MagickBooleanType regenerateProfile; /* re-run the microbenchmark in auto device selection mode */ + cl_device_type + type; - /* Random number generator seeds */ - unsigned int numGenerators; - float randNormalize; - cl_mem seeds; - SemaphoreInfo* seedsLock; + cl_platform_id + platform; - SemaphoreInfo* lock; + cl_program + program; - cl_command_queue commandQueues[MAX_COMMAND_QUEUES]; - ssize_t commandQueuesPos; - SemaphoreInfo* commandQueuesLock; + cl_uint + max_clock_frequency, + max_compute_units; + + cl_ulong + local_memory_size; + + double + score; + + MagickBooleanType + enabled; + + SemaphoreInfo + *lock; + + ssize_t + command_queues_index, + created_queues; + +#if MAGICKCORE_OPENCL_PROFILE_KERNELS + KernelProfileRecord + profileRecords[KERNEL_COUNT]; +#endif +}; + +struct _MagickCLEnv +{ + double + cpu_score; + + MagickBooleanType + enabled, + initialized; + + MagickCLDevice + *devices; + + MagickLibrary + *library; + + MagickThreadType + benchmark_thread_id; + + SemaphoreInfo + *lock; + + size_t + number_devices; }; #endif @@ -355,93 +400,38 @@ struct _MagickCLEnv { #define CLCharQuantumScale 72340172838076673.0f #endif -typedef enum { - AddNoiseKernel, - BlurRowKernel, - BlurColumnKernel, - CompositeKernel, - ComputeFunctionKernel, - ContrastKernel, - ContrastStretchKernel, - ConvolveKernel, - EqualizeKernel, - GrayScaleKernel, - HistogramKernel, - HullPass1Kernel, - HullPass2Kernel, - LocalContrastBlurRowKernel, - LocalContrastBlurApplyColumnKernel, - ModulateKernel, - MotionBlurKernel, - RandomNumberGeneratorKernel, - ResizeHorizontalKernel, - ResizeVerticalKernel, - RotationalBlurKernel, - UnsharpMaskBlurColumnKernel, - UnsharpMaskKernel, - WaveletDenoiseKernel, - KERNEL_COUNT -} ProfiledKernels; +extern MagickPrivate cl_command_queue + AcquireOpenCLCommandQueue(MagickCLDevice); -extern MagickPrivate cl_context - GetOpenCLContext(MagickCLEnv); +extern MagickPrivate cl_kernel + AcquireOpenCLKernel(MagickCLDevice,const char *); -extern MagickPrivate cl_kernel - AcquireOpenCLKernel(MagickCLEnv, MagickOpenCLProgram, const char*); +extern MagickPrivate MagickBooleanType + InitializeOpenCL(MagickCLEnv,ExceptionInfo *), + OpenCLThrowMagickException(MagickCLDevice,ExceptionInfo *, + const char *,const char *,const size_t,const ExceptionType,const char *, + const char *,...); -extern MagickPrivate cl_command_queue - AcquireOpenCLCommandQueue(MagickCLEnv); +extern MagickPrivate MagickCLDevice + GetOpenCLDevice(MagickCLEnv); -extern MagickPrivate MagickBooleanType - OpenCLThrowMagickException(ExceptionInfo *, - const char *,const char *,const size_t, - const ExceptionType,const char *,const char *,...), - RelinquishOpenCLCommandQueue(MagickCLEnv, cl_command_queue), - RelinquishOpenCLKernel(MagickCLEnv, cl_kernel); +extern MagickPrivate MagickCLEnv + GetCurrentOpenCLEnv(void); -extern MagickPrivate unsigned long - GetOpenCLDeviceLocalMemorySize(MagickCLEnv), - GetOpenCLDeviceMaxMemAllocSize(MagickCLEnv); - -extern MagickPrivate const char* - GetOpenCLCachedFilesDirectory(); - -extern MagickPrivate void - OpenCLLog(const char*), - UnlockRandSeedBuffer(MagickCLEnv); - -extern MagickPrivate cl_mem - GetAndLockRandSeedBuffer(MagickCLEnv); - -extern MagickPrivate unsigned int - GetNumRandGenerators(MagickCLEnv); - -extern MagickPrivate float - GetRandNormalize(MagickCLEnv); +extern MagickPrivate unsigned long + GetOpenCLDeviceLocalMemorySize(const MagickCLDevice); extern MagickPrivate void + DumpOpenCLProfileData(), OpenCLTerminus(), - RecordProfileData(MagickCLEnv,ProfiledKernels,cl_event); - -/* #define OPENCLLOG_ENABLED 1 */ -static inline void OpenCLLogException(const char* function, - const unsigned int line, - ExceptionInfo* exception) { -#ifdef OPENCLLOG_ENABLED - if (exception->severity!=0) { - char message[MagickPathExtent]; - /* dump the source into a file */ - (void) FormatLocaleString(message,MagickPathExtent,"%s:%d Exception(%d):%s " - ,function,line,exception->severity,exception->reason); - OpenCLLog(message); - } -#else - magick_unreferenced(function); - magick_unreferenced(line); - magick_unreferenced(exception); -#endif -} + RecordProfileData(MagickCLDevice,ProfiledKernels,cl_event), + RelinquishOpenCLCommandQueue(MagickCLDevice,cl_command_queue), + RelinquishOpenCLKernel(cl_kernel); +extern MagickPrivate cl_mem + GetAndLockRandSeedBuffer(MagickCLEnv,MagickCLDevice); +extern MagickPrivate void + UnlockRandSeedBuffer(MagickCLEnv); #if defined(__cplusplus) || defined(c_plusplus) } diff --git a/MagickCore/opencl.c b/MagickCore/opencl.c index f4df318ed4..0152259cc4 100644 --- a/MagickCore/opencl.c +++ b/MagickCore/opencl.c @@ -46,6 +46,7 @@ #include "MagickCore/color.h" #include "MagickCore/compare.h" #include "MagickCore/constitute.h" +#include "MagickCore/configure.h" #include "MagickCore/distort.h" #include "MagickCore/draw.h" #include "MagickCore/effect.h" @@ -79,1357 +80,466 @@ #include "MagickCore/semaphore.h" #include "MagickCore/statistic.h" #include "MagickCore/string_.h" +#include "MagickCore/string-private.h" #include "MagickCore/token.h" #include "MagickCore/utility.h" - -#ifdef MAGICKCORE_CLPERFMARKER -#include "CLPerfMarker.h" -#endif - +#include "MagickCore/utility-private.h" #if defined(MAGICKCORE_OPENCL_SUPPORT) +#ifndef MAGICKCORE_WINDOWS_SUPPORT +#include +#endif + #ifdef MAGICKCORE_HAVE_OPENCL_CL_H #define MAGICKCORE_OPENCL_MACOSX 1 #endif -#define NUM_CL_RAND_GENERATORS 1024 /* number of random number generators running in parallel */ -#define PROFILE_OCL_KERNELS 0 +/* + Define declarations. +*/ +#define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile" + +/* + Typedef declarations. +*/ +typedef struct +{ + long long freq; + long long clocks; + long long start; +} AccelerateTimer; typedef struct { - cl_ulong min; - cl_ulong max; - cl_ulong total; - cl_ulong count; -} KernelProfileRecord; + char + *name, + *version; -static const char *kernelNames[] = { + cl_uint + max_clock_frequency, + max_compute_units; + + double + score; +} MagickCLDeviceBenchmark; + +/* + Forward declarations. +*/ + +static MagickBooleanType + HasOpenCLDevices(MagickCLEnv,ExceptionInfo *), + LoadOpenCLLibrary(void); + +static MagickCLDevice + RelinquishMagickCLDevice(MagickCLDevice); + +static MagickCLEnv + RelinquishMagickCLEnv(MagickCLEnv); + +static void + BenchmarkOpenCLDevices(MagickCLEnv); + +extern const char + *accelerateKernels, *accelerateKernels2; + +/* + static declarations. +*/ +static const char *kernelNames[] = +{ "AddNoise", - "BlurRow", "BlurColumn", + "BlurRow", "Composite", - "ComputeFunction", "Contrast", "ContrastStretch", "Convolve", + "ConvolveOptimized", + "ComputeFunction", "Equalize", "GrayScale", "Histogram", "HullPass1", "HullPass2", - "LocalContrastBlurRow", "LocalContrastBlurApplyColumn", + "LocalContrastBlurRow", "Modulate", "MotionBlur", - "RandomNumberGenerator", "ResizeHorizontal", "ResizeVertical", "RotationalBlur", - "UnsharpMaskBlurColumn", "UnsharpMask", + "UnsharpMaskBlurColumn", "WaveletDenoise", - "NONE" }; - -KernelProfileRecord - profileRecords[KERNEL_COUNT]; - -typedef struct _AccelerateTimer { - long long _freq; - long long _clocks; - long long _start; -} AccelerateTimer; - -void startAccelerateTimer(AccelerateTimer* timer) { -#ifdef _WIN32 - QueryPerformanceCounter((LARGE_INTEGER*)&timer->_start); - - -#else - struct timeval s; - gettimeofday(&s, 0); - timer->_start = (long long)s.tv_sec * (long long)1.0E3 + (long long)s.tv_usec / (long long)1.0E3; -#endif -} - -void stopAccelerateTimer(AccelerateTimer* timer) { - long long n=0; -#ifdef _WIN32 - QueryPerformanceCounter((LARGE_INTEGER*)&(n)); -#else - struct timeval s; - gettimeofday(&s, 0); - n = (long long)s.tv_sec * (long long)1.0E3+ (long long)s.tv_usec / (long long)1.0E3; -#endif - n -= timer->_start; - timer->_start = 0; - timer->_clocks += n; -} - -void resetAccelerateTimer(AccelerateTimer* timer) { - timer->_clocks = 0; - timer->_start = 0; -} - -void initAccelerateTimer(AccelerateTimer* timer) { -#ifdef _WIN32 - QueryPerformanceFrequency((LARGE_INTEGER*)&timer->_freq); -#else - timer->_freq = (long long)1.0E3; -#endif - resetAccelerateTimer(timer); -} - -double readAccelerateTimer(AccelerateTimer* timer) { - return (double)timer->_clocks/(double)timer->_freq; + "NONE" }; -MagickPrivate void RecordProfileData(MagickCLEnv clEnv, ProfiledKernels kernel, cl_event event) +/* OpenCL library */ +MagickLibrary + *openCL_library; + +/* Default OpenCL environment */ +MagickCLEnv + default_CLEnv; +MagickThreadType + test_thread_id=0; +SemaphoreInfo + *default_CLEnv_Lock; + +/* Cached location of the OpenCL cache files */ +char + *cache_directory; +SemaphoreInfo + *cache_directory_lock; + +static inline MagickBooleanType IsSameOpenCLDevice(MagickCLDevice a, + MagickCLDevice b) { -#if PROFILE_OCL_KERNELS - cl_int status; - cl_ulong start = 0; - cl_ulong end = 0; - cl_ulong elapsed = 0; - clEnv->library->clWaitForEvents(1, &event); - status = clEnv->library->clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL); - status &= clEnv->library->clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL); - if (status == CL_SUCCESS) { - start /= 1000; // usecs - end /= 1000; // usecs - elapsed = end - start; - /* we can use the commandQueuesLock to make the code below thread safe */ - LockSemaphoreInfo(clEnv->commandQueuesLock); - if ((elapsed < profileRecords[kernel].min) || (profileRecords[kernel].count == 0)) - profileRecords[kernel].min = elapsed; - if (elapsed > profileRecords[kernel].max) - profileRecords[kernel].max = elapsed; - profileRecords[kernel].total += elapsed; - profileRecords[kernel].count += 1; - UnlockSemaphoreInfo(clEnv->commandQueuesLock); - } -#else - magick_unreferenced(clEnv); - magick_unreferenced(kernel); - magick_unreferenced(event); -#endif + if ((LocaleCompare(a->name,b->name) == 0) && + (LocaleCompare(a->version,b->version) == 0) && + (a->max_clock_frequency == b->max_clock_frequency) && + (a->max_compute_units == b->max_compute_units)) + return(MagickTrue); + + return(MagickFalse); } -void DumpProfileData() +static inline MagickBooleanType IsBenchmarkedOpenCLDevice(MagickCLDevice a, + MagickCLDeviceBenchmark *b) { -#if PROFILE_OCL_KERNELS - int i; + if ((LocaleCompare(a->name,b->name) == 0) && + (LocaleCompare(a->version,b->version) == 0) && + (a->max_clock_frequency == b->max_clock_frequency) && + (a->max_compute_units == b->max_compute_units)) + return(MagickTrue); - OpenCLLog("===================================================="); - - // Write out the device info to the profile - if (0 == 1) - { - MagickCLEnv clEnv; - char buff[2048]; - cl_int status; - - clEnv = GetDefaultOpenCLEnv(); - - status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_VENDOR, 2048, buff, NULL); - OpenCLLog(buff); - - status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_NAME, 2048, buff, NULL); - OpenCLLog(buff); - - status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DRIVER_VERSION, 2048, buff, NULL); - OpenCLLog(buff); - } - - OpenCLLog("===================================================="); - OpenCLLog(" ave\tcalls \tmin -> max"); - OpenCLLog(" ---\t----- \t----------"); - for (i = 0; i < KERNEL_COUNT; ++i) { - char buf[4096]; - char indent[160]; - strcpy(indent, " "); - strncpy(indent, kernelNames[i], min(strlen(kernelNames[i]), strlen(indent) - 1)); - sprintf(buf, "%s%d\t(%d calls) \t%d -> %d", indent, profileRecords[i].count > 0 ? (profileRecords[i].total / profileRecords[i].count) : 0, profileRecords[i].count, profileRecords[i].min, profileRecords[i].max); - //printf("%s%d\t(%d calls) \t%d -> %d\n", indent, profileRecords[i].count > 0 ? (profileRecords[i].total / profileRecords[i].count) : 0, profileRecords[i].count, profileRecords[i].min, profileRecords[i].max); - OpenCLLog(buf); - } - OpenCLLog("===================================================="); -#endif + return(MagickFalse); } -/* - * - * Dynamic library loading functions - * - */ -#ifdef MAGICKCORE_WINDOWS_SUPPORT -#else -#include -#endif - -// dynamically load a library. returns NULL on failure -void *OsLibraryLoad(const char *libraryName) +static inline void RelinquishMagickCLDevices(MagickCLEnv clEnv) { -#ifdef MAGICKCORE_WINDOWS_SUPPORT - return (void *)LoadLibraryA(libraryName); -#else - return (void *)dlopen(libraryName, RTLD_NOW); -#endif -} + size_t + i; -// get a function pointer from a loaded library. returns NULL on failure. -void *OsLibraryGetFunctionAddress(void *library, const char *functionName) -{ -#ifdef MAGICKCORE_WINDOWS_SUPPORT - if (!library || !functionName) + if (clEnv->devices != (MagickCLDevice *) NULL) { - return NULL; + for (i = 0; i < clEnv->number_devices; i++) + clEnv->devices[i]=RelinquishMagickCLDevice(clEnv->devices[i]); + clEnv->devices=(MagickCLDevice *) RelinquishMagickMemory(clEnv->devices); } - return (void *) GetProcAddress( (HMODULE)library, functionName); -#else - if (!library || !functionName) - { - return NULL; - } - return (void *)dlsym(library, functionName); -#endif + clEnv->number_devices=0; } -// unload a library. -void OsLibraryUnload(void *library) +static inline MagickBooleanType MagickCreateDirectory(const char *path) { + int + status; + #ifdef MAGICKCORE_WINDOWS_SUPPORT - FreeLibrary( (HMODULE)library); + status=mkdir(path); #else - dlclose(library); + status=mkdir(path, 0777); +#endif + return(status == 0 ? MagickTrue : MagickFalse); +} + +static inline void InitAccelerateTimer(AccelerateTimer *timer) +{ +#ifdef _WIN32 + QueryPerformanceFrequency((LARGE_INTEGER*)&timer->freq); +#else + timer->freq=(long long)1.0E3; +#endif + timer->clocks=0; + timer->start=0; +} + +static inline double ReadAccelerateTimer(AccelerateTimer *timer) +{ + return (double)timer->clocks/(double)timer->freq; +} + +static inline void StartAccelerateTimer(AccelerateTimer* timer) +{ +#ifdef _WIN32 + QueryPerformanceCounter((LARGE_INTEGER*)&timer->start); +#else + struct timeval + s; + gettimeofday(&s,0); + timer->start=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/ + (long long)1.0E3; #endif } - -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -+ A c q u i r e M a g i c k O p e n C L E n v % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% -% AcquireMagickOpenCLEnv() allocates the MagickCLEnv structure -% -*/ - -MagickExport MagickCLEnv AcquireMagickOpenCLEnv(void) +static inline void StopAccelerateTimer(AccelerateTimer *timer) { - MagickCLEnv clEnv; - clEnv = (MagickCLEnv) AcquireMagickMemory(sizeof(struct _MagickCLEnv)); - if (clEnv != NULL) - { - memset(clEnv, 0, sizeof(struct _MagickCLEnv)); - clEnv->commandQueuesPos=-1; - ActivateSemaphoreInfo(&clEnv->lock); - ActivateSemaphoreInfo(&clEnv->commandQueuesLock); - } - return clEnv; -} + long long + n; - -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -+ R e l i n q u i s h M a g i c k O p e n C L E n v % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% -% RelinquishMagickOpenCLEnv() destroy the MagickCLEnv structure -% -% The format of the RelinquishMagickOpenCLEnv method is: -% -% MagickBooleanType RelinquishMagickOpenCLEnv(MagickCLEnv clEnv) -% -% A description of each parameter follows: -% -% o clEnv: MagickCLEnv structure to destroy -% -*/ - -MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(MagickCLEnv clEnv) -{ - if (clEnv != (MagickCLEnv) NULL) - { - while (clEnv->commandQueuesPos >= 0) - { - clEnv->library->clReleaseCommandQueue( - clEnv->commandQueues[clEnv->commandQueuesPos--]); - } - RelinquishSemaphoreInfo(&clEnv->lock); - RelinquishSemaphoreInfo(&clEnv->commandQueuesLock); - RelinquishMagickMemory(clEnv); - return MagickTrue; - } - return MagickFalse; -} - - -/* -* Default OpenCL environment -*/ -MagickCLEnv defaultCLEnv; -SemaphoreInfo* defaultCLEnvLock; - -/* -* OpenCL library -*/ -MagickLibrary * OpenCLLib; -SemaphoreInfo* OpenCLLibLock; - - -static MagickBooleanType bindOpenCLFunctions(void* library) -{ -#ifdef MAGICKCORE_OPENCL_MACOSX -#define BIND(X) OpenCLLib->X= &X; + n=0; +#ifdef _WIN32 + QueryPerformanceCounter((LARGE_INTEGER*)&(n)); #else -#define BIND(X)\ - if ((OpenCLLib->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(library,#X)) == NULL)\ - return MagickFalse; + struct timeval + s; + gettimeofday(&s,0); + n=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/ + (long long)1.0E3; #endif - - BIND(clGetPlatformIDs); - BIND(clGetPlatformInfo); - - BIND(clGetDeviceIDs); - BIND(clGetDeviceInfo); - - BIND(clCreateContext); - - BIND(clCreateBuffer); - BIND(clReleaseMemObject); - - BIND(clCreateProgramWithSource); - BIND(clCreateProgramWithBinary); - BIND(clBuildProgram); - BIND(clGetProgramInfo); - BIND(clGetProgramBuildInfo); - - BIND(clCreateKernel); - BIND(clReleaseKernel); - BIND(clSetKernelArg); - - BIND(clFlush); - BIND(clFinish); - - BIND(clEnqueueNDRangeKernel); - BIND(clEnqueueReadBuffer); - BIND(clEnqueueMapBuffer); - BIND(clEnqueueUnmapMemObject); - - BIND(clCreateCommandQueue); - BIND(clReleaseCommandQueue); - - BIND(clGetEventProfilingInfo); - BIND(clWaitForEvents); - BIND(clReleaseEvent); - - return MagickTrue; + n-=timer->start; + timer->start=0; + timer->clocks+=n; } -MagickLibrary * GetOpenCLLib() +static const char *GetOpenCLCacheDirectory() { - if (OpenCLLib == NULL) - { - if (OpenCLLibLock == NULL) + if (cache_directory != (char *) NULL) + return(cache_directory); + + if (cache_directory_lock == (SemaphoreInfo *) NULL) + ActivateSemaphoreInfo(&cache_directory_lock); + LockSemaphoreInfo(cache_directory_lock); + if (cache_directory == (char *) NULL) { - ActivateSemaphoreInfo(&OpenCLLibLock); - } + char + *home, + path[MagickPathExtent], + *temp; - LockSemaphoreInfo(OpenCLLibLock); + MagickBooleanType + status; - OpenCLLib = (MagickLibrary *) AcquireMagickMemory (sizeof (MagickLibrary)); + struct stat + attributes; - if (OpenCLLib != NULL) - { - MagickBooleanType status = MagickFalse; - void * library = NULL; + home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR"); + if (home == (char *) NULL) + { + home=GetEnvironmentValue("XDG_CACHE_HOME"); + if (home == (char *) NULL) + home=GetEnvironmentValue("LOCALAPPDATA"); + if (home == (char *) NULL) + home=GetEnvironmentValue("APPDATA"); + if (home == (char *) NULL) + home=GetEnvironmentValue("USERPROFILE"); + } -#ifdef MAGICKCORE_OPENCL_MACOSX - status = bindOpenCLFunctions(library); -#else + if (home != (char *) NULL) + { + /* first check if $HOME exists */ + (void) FormatLocaleString(path,MagickPathExtent,"%s",home); + status=GetPathAttributes(path,&attributes); + if (status == MagickFalse) + status=MagickCreateDirectory(path); - memset(OpenCLLib, 0, sizeof(MagickLibrary)); -#ifdef MAGICKCORE_WINDOWS_SUPPORT - library = OsLibraryLoad("OpenCL.dll"); -#else - library = OsLibraryLoad("libOpenCL.so"); -#endif - if (library) - status = bindOpenCLFunctions(library); + /* first check if $HOME/ImageMagick exists */ + if (status != MagickFalse) + { + (void) FormatLocaleString(path,MagickPathExtent, + "%s%sImageMagick",home,DirectorySeparator); - if (status==MagickTrue) - OpenCLLib->base=library; + status=GetPathAttributes(path,&attributes); + if (status == MagickFalse) + status=MagickCreateDirectory(path); + } + + if (status != MagickFalse) + { + temp=(char*)AcquireMagickMemory(strlen(path)+1); + CopyMagickString(temp,path,strlen(path)+1); + } + home=DestroyString(home); + } else - OpenCLLib=(MagickLibrary *)RelinquishMagickMemory(OpenCLLib); -#endif + { + home=GetEnvironmentValue("HOME"); + if (home != (char *) NULL) + { + /* first check if $HOME/.cache exists */ + (void) FormatLocaleString(path,MagickPathExtent,"%s%s.cache", + home,DirectorySeparator); + status=GetPathAttributes(path,&attributes); + if (status == MagickFalse) + status=MagickCreateDirectory(path); + + /* first check if $HOME/.cache/ImageMagick exists */ + if (status != MagickFalse) + { + (void) FormatLocaleString(path,MagickPathExtent, + "%s%s.cache%sImageMagick",home,DirectorySeparator, + DirectorySeparator); + status=GetPathAttributes(path,&attributes); + if (status == MagickFalse) + status=MagickCreateDirectory(path); + } + + if (status != MagickFalse) + { + temp=(char*)AcquireMagickMemory(strlen(path)+1); + CopyMagickString(temp,path,strlen(path)+1); + } + home=DestroyString(home); + } + } + cache_directory=temp; } - - UnlockSemaphoreInfo(OpenCLLibLock); - } - - - return OpenCLLib; + UnlockSemaphoreInfo(cache_directory_lock); + return(cache_directory); } - -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -+ G e t D e f a u l t O p e n C L E n v % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% -% GetDefaultOpenCLEnv() returns the default OpenCL env -% -% The format of the GetDefaultOpenCLEnv method is: -% -% MagickCLEnv GetDefaultOpenCLEnv() -% -% A description of each parameter follows: -% -% o exception: return any errors or warnings. -% -*/ - -MagickExport MagickCLEnv GetDefaultOpenCLEnv(void) +static void SelectOpenCLDevice(MagickCLEnv clEnv,cl_device_type type) { - if (defaultCLEnv == NULL) + MagickCLDevice + device; + + size_t + i, + j; + + for (i = 0; i < clEnv->number_devices; i++) + clEnv->devices[i]->enabled=MagickFalse; + + for (i = 0; i < clEnv->number_devices; i++) { - if (defaultCLEnvLock == NULL) + device=clEnv->devices[i]; + if (device->type != type) + continue; + + device->enabled=MagickTrue; + for (j = i+1; j < clEnv->number_devices; j++) { - ActivateSemaphoreInfo(&defaultCLEnvLock); + MagickCLDevice + other_device; + + other_device=clEnv->devices[j]; + if (IsSameOpenCLDevice(device,other_device)) + other_device->enabled=MagickTrue; } - LockSemaphoreInfo(defaultCLEnvLock); - if (defaultCLEnv == NULL) - defaultCLEnv = AcquireMagickOpenCLEnv(); - UnlockSemaphoreInfo(defaultCLEnvLock); } - return defaultCLEnv; } -static void LockDefaultOpenCLEnv() { - if (defaultCLEnvLock == NULL) - { - ActivateSemaphoreInfo(&defaultCLEnvLock); - } - LockSemaphoreInfo(defaultCLEnvLock); -} - -static void UnlockDefaultOpenCLEnv() { - if (defaultCLEnvLock == NULL) - { - ActivateSemaphoreInfo(&defaultCLEnvLock); - } - else - UnlockSemaphoreInfo(defaultCLEnvLock); -} - - -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -+ S e t D e f a u l t O p e n C L E n v % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% -% SetDefaultOpenCLEnv() sets the new OpenCL environment as default -% and returns the old OpenCL environment -% -% The format of the SetDefaultOpenCLEnv() method is: -% -% MagickCLEnv SetDefaultOpenCLEnv(MagickCLEnv clEnv) -% -% A description of each parameter follows: -% -% o clEnv: the new default OpenCL environment. -% -*/ -MagickExport MagickCLEnv SetDefaultOpenCLEnv(MagickCLEnv clEnv) +static unsigned int StringSignature(const char* string) { - MagickCLEnv oldEnv; - LockDefaultOpenCLEnv(); - oldEnv = defaultCLEnv; - defaultCLEnv = clEnv; - UnlockDefaultOpenCLEnv(); - return oldEnv; -} + unsigned int + n, + i, + j, + signature, + stringLength; - - -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -+ S e t M a g i c k O p e n C L E n v P a r a m % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% -% SetMagickOpenCLEnvParam() sets the parameters in the OpenCL environment -% -% The format of the SetMagickOpenCLEnvParam() method is: -% -% MagickBooleanType SetMagickOpenCLEnvParam(MagickCLEnv clEnv, -% MagickOpenCLEnvParam param, size_t dataSize, void* data, -% ExceptionInfo* exception) -% -% A description of each parameter follows: -% -% o clEnv: the OpenCL environment. -% -% o param: the parameter to be set. -% -% o dataSize: the data size of the parameter value. -% -% o data: the pointer to the new parameter value -% -% o exception: return any errors or warnings -% -*/ - -static MagickBooleanType SetMagickOpenCLEnvParamInternal(MagickCLEnv clEnv, MagickOpenCLEnvParam param - , size_t dataSize, void* data, ExceptionInfo* exception) -{ - MagickBooleanType status = MagickFalse; - - if (clEnv == NULL - || data == NULL) - goto cleanup; - - switch(param) - { - case MAGICK_OPENCL_ENV_PARAM_DEVICE: - if (dataSize != sizeof(clEnv->device)) - goto cleanup; - clEnv->device = *((cl_device_id*)data); - clEnv->OpenCLInitialized = MagickFalse; - status = MagickTrue; - break; - - case MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED: - if (dataSize != sizeof(clEnv->OpenCLDisabled)) - goto cleanup; - clEnv->OpenCLDisabled = *((MagickBooleanType*)data); - clEnv->OpenCLInitialized = MagickFalse; - status = MagickTrue; - break; - - case MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED: - (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "SetMagickOpenCLEnvParm cannot modify the OpenCL initialization state.", "'%s'", "."); - break; - - case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED: - if (dataSize != sizeof(clEnv->disableProgramCache)) - goto cleanup; - clEnv->disableProgramCache = *((MagickBooleanType*)data); - clEnv->OpenCLInitialized = MagickFalse; - status = MagickTrue; - break; - - case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE: - if (dataSize != sizeof(clEnv->regenerateProfile)) - goto cleanup; - clEnv->regenerateProfile = *((MagickBooleanType*)data); - clEnv->OpenCLInitialized = MagickFalse; - status = MagickTrue; - break; - - default: - goto cleanup; - }; - -cleanup: - return status; -} - -MagickExport - MagickBooleanType SetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param - , size_t dataSize, void* data, ExceptionInfo* exception) { - MagickBooleanType status = MagickFalse; - if (clEnv!=NULL) { - LockSemaphoreInfo(clEnv->lock); - status = SetMagickOpenCLEnvParamInternal(clEnv,param,dataSize,data,exception); - UnlockSemaphoreInfo(clEnv->lock); - } - return status; -} - -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -+ G e t M a g i c k O p e n C L E n v P a r a m % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% -% GetMagickOpenCLEnvParam() gets the parameters in the OpenCL environment -% -% The format of the GetMagickOpenCLEnvParam() method is: -% -% MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv, -% MagickOpenCLEnvParam param, size_t dataSize, void* data, -% ExceptionInfo* exception) -% -% A description of each parameter follows: -% -% o clEnv: the OpenCL environment. -% -% o param: the parameter to be returned. -% -% o dataSize: the data size of the parameter value. -% -% o data: the location where the returned parameter value will be stored -% -% o exception: return any errors or warnings -% -*/ - -MagickExport - MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param - , size_t dataSize, void* data, ExceptionInfo* exception) -{ - MagickBooleanType - status; - - magick_unreferenced(exception); - - status = MagickFalse; - - if (clEnv == NULL - || data == NULL) - goto cleanup; - - switch(param) - { - case MAGICK_OPENCL_ENV_PARAM_DEVICE: - if (dataSize != sizeof(cl_device_id)) - goto cleanup; - *((cl_device_id*)data) = clEnv->device; - status = MagickTrue; - break; - - case MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED: - if (dataSize != sizeof(clEnv->OpenCLDisabled)) - goto cleanup; - *((MagickBooleanType*)data) = clEnv->OpenCLDisabled; - status = MagickTrue; - break; - - case MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED: - if (dataSize != sizeof(clEnv->OpenCLDisabled)) - goto cleanup; - *((MagickBooleanType*)data) = clEnv->OpenCLInitialized; - status = MagickTrue; - break; - - case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED: - if (dataSize != sizeof(clEnv->disableProgramCache)) - goto cleanup; - *((MagickBooleanType*)data) = clEnv->disableProgramCache; - status = MagickTrue; - break; - - case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE: - if (dataSize != sizeof(clEnv->regenerateProfile)) - goto cleanup; - *((MagickBooleanType*)data) = clEnv->regenerateProfile; - status = MagickTrue; - break; - - default: - goto cleanup; - }; - -cleanup: - return status; -} - - -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -+ G e t O p e n C L C o n t e x t % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% -% GetOpenCLContext() returns the OpenCL context -% -% The format of the GetOpenCLContext() method is: -% -% cl_context GetOpenCLContext(MagickCLEnv clEnv) -% -% A description of each parameter follows: -% -% o clEnv: OpenCL environment -% -*/ - -MagickPrivate -cl_context GetOpenCLContext(MagickCLEnv clEnv) { - if (clEnv == NULL) - return NULL; - else - return clEnv->context; -} - -static char* getBinaryCLProgramName(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature) -{ - char* name; - char* ptr; - char path[MagickPathExtent]; - char deviceName[MagickPathExtent]; - const char* prefix = "magick_opencl"; - clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_NAME, MagickPathExtent, deviceName, NULL); - ptr=deviceName; - /* strip out illegal characters for file names */ - while (*ptr != '\0') - { - if ( *ptr == ' ' || *ptr == '\\' || *ptr == '/' || *ptr == ':' || *ptr == '*' - || *ptr == '?' || *ptr == '"' || *ptr == '<' || *ptr == '>' || *ptr == '|') - { - *ptr = '_'; - } - ptr++; - } - (void) FormatLocaleString(path,MagickPathExtent,"%s%s%s_%s_%02d_%08x_%.20g.bin", - GetOpenCLCachedFilesDirectory(),DirectorySeparator,prefix,deviceName, - (unsigned int) prog,signature,(double) sizeof(char*)*8); - name = (char*)AcquireMagickMemory(strlen(path)+1); - CopyMagickString(name,path,strlen(path)+1); - return name; -} - -static MagickBooleanType saveBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature, ExceptionInfo* exception) -{ - MagickBooleanType saveSuccessful; - cl_int clStatus; - size_t binaryProgramSize; - unsigned char* binaryProgram; - char* binaryFileName; - FILE* fileHandle; - -#ifdef MAGICKCORE_CLPERFMARKER - clBeginPerfMarkerAMD(__FUNCTION__,""); -#endif - - binaryProgram = NULL; - binaryFileName = NULL; - fileHandle = NULL; - saveSuccessful = MagickFalse; - - clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binaryProgramSize, NULL); - if (clStatus != CL_SUCCESS) - { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", "."); - goto cleanup; - } - - binaryProgram = (unsigned char*) AcquireMagickMemory(binaryProgramSize); - clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARIES, sizeof(char*), &binaryProgram, NULL); - if (clStatus != CL_SUCCESS) - { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", "."); - goto cleanup; - } - - binaryFileName = getBinaryCLProgramName(clEnv, prog, signature); - fileHandle = fopen(binaryFileName, "wb"); - if (fileHandle != NULL) - { - fwrite(binaryProgram, sizeof(char), binaryProgramSize, fileHandle); - saveSuccessful = MagickTrue; - } - else - { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "Saving binary kernel failed.", "'%s'", "."); - } - -cleanup: - if (fileHandle != NULL) - fclose(fileHandle); - if (binaryProgram != NULL) - RelinquishMagickMemory(binaryProgram); - if (binaryFileName != NULL) - free(binaryFileName); - -#ifdef MAGICKCORE_CLPERFMARKER - clEndPerfMarkerAMD(); -#endif - - return saveSuccessful; -} - -static MagickBooleanType loadBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature) -{ - MagickBooleanType loadSuccessful; - unsigned char* binaryProgram; - char* binaryFileName; - FILE* fileHandle; - -#ifdef MAGICKCORE_CLPERFMARKER - clBeginPerfMarkerAMD(__FUNCTION__,""); -#endif - - binaryProgram = NULL; - binaryFileName = NULL; - fileHandle = NULL; - loadSuccessful = MagickFalse; - - binaryFileName = getBinaryCLProgramName(clEnv, prog, signature); - fileHandle = fopen(binaryFileName, "rb"); - if (fileHandle != NULL) - { - int b_error; - size_t length; - cl_int clStatus; - cl_int clBinaryStatus; - - b_error = 0 ; - length = 0; - b_error |= fseek( fileHandle, 0, SEEK_END ) < 0; - b_error |= ( length = ftell( fileHandle ) ) <= 0; - b_error |= fseek( fileHandle, 0, SEEK_SET ) < 0; - if( b_error ) - goto cleanup; - - binaryProgram = (unsigned char*)AcquireMagickMemory(length); - if (binaryProgram == NULL) - goto cleanup; - - memset(binaryProgram, 0, length); - b_error |= fread(binaryProgram, 1, length, fileHandle) != length; - - clEnv->programs[prog] = clEnv->library->clCreateProgramWithBinary(clEnv->context, 1, &clEnv->device, &length, (const unsigned char**)&binaryProgram, &clBinaryStatus, &clStatus); - if (clStatus != CL_SUCCESS - || clBinaryStatus != CL_SUCCESS) - goto cleanup; - - loadSuccessful = MagickTrue; - } - -cleanup: - if (fileHandle != NULL) - fclose(fileHandle); - if (binaryFileName != NULL) - free(binaryFileName); - if (binaryProgram != NULL) - RelinquishMagickMemory(binaryProgram); - -#ifdef MAGICKCORE_CLPERFMARKER - clEndPerfMarkerAMD(); -#endif - - return loadSuccessful; -} - -static unsigned int stringSignature(const char* string) -{ - unsigned int stringLength; - unsigned int n,i,j; - unsigned int signature; union { const char* s; const unsigned int* u; - }p; + } p; -#ifdef MAGICKCORE_CLPERFMARKER - clBeginPerfMarkerAMD(__FUNCTION__,""); -#endif - - stringLength = (unsigned int) strlen(string); - signature = stringLength; - n = stringLength/sizeof(unsigned int); - p.s = string; + stringLength=(unsigned int) strlen(string); + signature=stringLength; + n=stringLength/sizeof(unsigned int); + p.s=string; for (i = 0; i < n; i++) - { signature^=p.u[i]; - } if (n * sizeof(unsigned int) != stringLength) - { - char padded[4]; - j = n * sizeof(unsigned int); - for (i = 0; i < 4; i++,j++) { - if (j < stringLength) - padded[i] = p.s[j]; - else - padded[i] = 0; + char + padded[4]; + + j=n*sizeof(unsigned int); + for (i = 0; i < 4; i++, j++) + { + if (j < stringLength) + padded[i]=p.s[j]; + else + padded[i]=0; + } + p.s=padded; + signature^=p.u[0]; } - p.s = padded; - signature^=p.u[0]; - } - -#ifdef MAGICKCORE_CLPERFMARKER - clEndPerfMarkerAMD(); -#endif - - return signature; + return(signature); } -/* OpenCL kernels for accelerate.c */ -extern const char *accelerateKernels, *accelerateKernels2; - -static MagickBooleanType CompileOpenCLKernels(MagickCLEnv clEnv, ExceptionInfo* exception) -{ - MagickBooleanType status = MagickFalse; - cl_int clStatus; - unsigned int i; - char* accelerateKernelsBuffer = NULL; - - /* The index of the program strings in this array has to match the value of the enum MagickOpenCLProgram */ - const char* MagickOpenCLProgramStrings[MAGICK_OPENCL_NUM_PROGRAMS]; - - char options[MagickPathExtent]; - unsigned int optionsSignature; - -#ifdef MAGICKCORE_CLPERFMARKER - clBeginPerfMarkerAMD(__FUNCTION__,""); -#endif - - /* Get additional options */ - (void) FormatLocaleString(options, MagickPathExtent, CLOptions, (float)QuantumRange, - (float)QuantumScale, (float)CLCharQuantumScale, (float)MagickEpsilon, (float)MagickPI, (unsigned int)MaxMap, (unsigned int)MAGICKCORE_QUANTUM_DEPTH); - - /* - if (getenv("MAGICK_OCL_DEF")) - { - strcat(options," "); - strcat(options,getenv("MAGICK_OCL_DEF")); - } - */ - - /* - if (getenv("MAGICK_OCL_BUILD")) - printf("options: %s\n", options); - */ - - optionsSignature = stringSignature(options); - - /* get all the OpenCL program strings here */ - accelerateKernelsBuffer = (char*) AcquireMagickMemory(strlen(accelerateKernels)+strlen(accelerateKernels2)+1); - sprintf(accelerateKernelsBuffer,"%s%s",accelerateKernels,accelerateKernels2); - MagickOpenCLProgramStrings[MAGICK_OPENCL_ACCELERATE] = accelerateKernelsBuffer; - - for (i = 0; i < MAGICK_OPENCL_NUM_PROGRAMS; i++) - { - MagickBooleanType loadSuccessful = MagickFalse; - unsigned int programSignature = stringSignature(MagickOpenCLProgramStrings[i]) ^ optionsSignature; - - /* try to load the binary first */ - if (clEnv->disableProgramCache != MagickTrue - && !getenv("MAGICK_OCL_REC")) - loadSuccessful = loadBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature); - - if (loadSuccessful == MagickFalse) - { - /* Binary CL program unavailable, compile the program from source */ - size_t programLength = strlen(MagickOpenCLProgramStrings[i]); - clEnv->programs[i] = clEnv->library->clCreateProgramWithSource(clEnv->context, 1, &(MagickOpenCLProgramStrings[i]), &programLength, &clStatus); - if (clStatus!=CL_SUCCESS) - { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "clCreateProgramWithSource failed.", "(%d)", (int)clStatus); - - goto cleanup; - } - } - - clStatus = clEnv->library->clBuildProgram(clEnv->programs[i], 1, &clEnv->device, options, NULL, NULL); - if (clStatus!=CL_SUCCESS) - { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "clBuildProgram failed.", "(%d)", (int)clStatus); - - if (loadSuccessful == MagickFalse) - { - char path[MagickPathExtent]; - FILE* fileHandle; - - /* dump the source into a file */ - (void) FormatLocaleString(path,MagickPathExtent,"%s%s%s" - ,GetOpenCLCachedFilesDirectory() - ,DirectorySeparator,"magick_badcl.cl"); - fileHandle = fopen(path, "wb"); - if (fileHandle != NULL) - { - fwrite(MagickOpenCLProgramStrings[i], sizeof(char), strlen(MagickOpenCLProgramStrings[i]), fileHandle); - fclose(fileHandle); - } - - /* dump the build log */ - { - char* log; - size_t logSize; - clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize); - log = (char*)AcquireMagickMemory(logSize); - clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, logSize, log, &logSize); - - (void) FormatLocaleString(path,MagickPathExtent,"%s%s%s" - ,GetOpenCLCachedFilesDirectory() - ,DirectorySeparator,"magick_badcl_build.log"); - fileHandle = fopen(path, "wb"); - if (fileHandle != NULL) - { - const char* buildOptionsTitle = "build options: "; - fwrite(buildOptionsTitle, sizeof(char), strlen(buildOptionsTitle), fileHandle); - fwrite(options, sizeof(char), strlen(options), fileHandle); - fwrite("\n",sizeof(char), 1, fileHandle); - fwrite(log, sizeof(char), logSize, fileHandle); - fclose(fileHandle); - } - RelinquishMagickMemory(log); - } - } - goto cleanup; - } - - if (loadSuccessful == MagickFalse) - { - /* Save the binary to a file to avoid re-compilation of the kernels in the future */ - saveBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature, exception); - } - - } - status = MagickTrue; - -cleanup: - - if (accelerateKernelsBuffer!=NULL) RelinquishMagickMemory(accelerateKernelsBuffer); - -#ifdef MAGICKCORE_CLPERFMARKER - clEndPerfMarkerAMD(); -#endif - - return status; -} - -static MagickBooleanType InitOpenCLPlatformDevice(MagickCLEnv clEnv, ExceptionInfo* exception) { - int i,j; - cl_int status; - cl_uint numPlatforms = 0; - cl_platform_id *platforms = NULL; - char* MAGICK_OCL_DEVICE = NULL; - MagickBooleanType OpenCLAvailable = MagickFalse; - -#ifdef MAGICKCORE_CLPERFMARKER - clBeginPerfMarkerAMD(__FUNCTION__,""); -#endif - - /* check if there's an environment variable overriding the device selection */ - MAGICK_OCL_DEVICE = getenv("MAGICK_OCL_DEVICE"); - if (MAGICK_OCL_DEVICE != NULL) - { - if (strcmp(MAGICK_OCL_DEVICE, "CPU") == 0) - { - clEnv->deviceType = CL_DEVICE_TYPE_CPU; - } - else if (strcmp(MAGICK_OCL_DEVICE, "GPU") == 0) - { - clEnv->deviceType = CL_DEVICE_TYPE_GPU; - } - else if (strcmp(MAGICK_OCL_DEVICE, "OFF") == 0) - { - /* OpenCL disabled */ - goto cleanup; - } - } - else if (clEnv->deviceType == 0) { - clEnv->deviceType = CL_DEVICE_TYPE_ALL; - } - - if (clEnv->device != NULL) - { - status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &clEnv->platform, NULL); - if (status != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "Failed to get OpenCL platform from the selected device.", "(%d)", status); - } - goto cleanup; - } - else if (clEnv->platform != NULL) - { - numPlatforms = 1; - platforms = (cl_platform_id *) AcquireMagickMemory(numPlatforms * sizeof(cl_platform_id)); - if (platforms == (cl_platform_id *) NULL) - { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError, - "AcquireMagickMemory failed.","."); - goto cleanup; - } - platforms[0] = clEnv->platform; - } - else - { - clEnv->device = NULL; - - /* Get the number of OpenCL platforms available */ - status = clEnv->library->clGetPlatformIDs(0, NULL, &numPlatforms); - if (status != CL_SUCCESS) - { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "clGetplatformIDs failed.", "(%d)", status); - goto cleanup; - } - - /* No OpenCL available, just leave */ - if (numPlatforms == 0) { - goto cleanup; - } - - platforms = (cl_platform_id *) AcquireMagickMemory(numPlatforms * sizeof(cl_platform_id)); - if (platforms == (cl_platform_id *) NULL) - { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError, - "AcquireMagickMemory failed.","."); - goto cleanup; - } - - status = clEnv->library->clGetPlatformIDs(numPlatforms, platforms, NULL); - if (status != CL_SUCCESS) - { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "clGetPlatformIDs failed.", "(%d)", status); - goto cleanup; - } - } - - /* Device selection */ - clEnv->device = NULL; - for (j = 0; j < 2; j++) - { - - cl_device_type deviceType; - if (clEnv->deviceType == CL_DEVICE_TYPE_ALL) - { - if (j == 0) - deviceType = CL_DEVICE_TYPE_GPU; - else - deviceType = CL_DEVICE_TYPE_CPU; - } - else if (j == 1) - { - break; - } - else - deviceType = clEnv->deviceType; - - for (i = 0; i < numPlatforms; i++) - { - char version[MagickPathExtent]; - cl_uint numDevices; - status = clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_VERSION, MagickPathExtent, version, NULL); - if (status != CL_SUCCESS) - { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "clGetPlatformInfo failed.", "(%d)", status); - goto cleanup; - } - if (strncmp(version,"OpenCL 1.0 ",11) == 0) - continue; - status = clEnv->library->clGetDeviceIDs(platforms[i], deviceType, 1, &(clEnv->device), &numDevices); - if (status != CL_SUCCESS) - { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "clGetDeviceIDs failed.", "(%d)", status); - goto cleanup; - } - if (clEnv->device != NULL) - { - clEnv->platform = platforms[i]; - goto cleanup; - } - } - } - -cleanup: - if (platforms!=NULL) - RelinquishMagickMemory(platforms); - - OpenCLAvailable = (clEnv->platform!=NULL - && clEnv->device!=NULL)?MagickTrue:MagickFalse; - -#ifdef MAGICKCORE_CLPERFMARKER - clEndPerfMarkerAMD(); -#endif - - return OpenCLAvailable; -} - -static MagickBooleanType EnableOpenCLInternal(MagickCLEnv clEnv) { - if (clEnv->OpenCLInitialized != MagickFalse - && clEnv->platform != NULL - && clEnv->device != NULL) { - clEnv->OpenCLDisabled = MagickFalse; - return MagickTrue; - } - clEnv->OpenCLDisabled = MagickTrue; - return MagickFalse; -} - - -static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception); /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % % % % % -+ I n i t O p e n C L E n v % +% A c q u i r e M a g i c k C L D e v i c e % % % % % % % %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % -% InitOpenCLEnv() initialize the OpenCL environment +% AcquireMagickCLDevice() acquires an OpenCL device % -% The format of the RelinquishMagickOpenCLEnv method is: +% The format of the AcquireMagickCLDevice method is: % -% MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception) -% -% A description of each parameter follows: -% -% o clEnv: OpenCL environment structure -% -% o exception: return any errors or warnings. +% MagickCLDevice AcquireMagickCLDevice() % */ -MagickExport -MagickBooleanType InitOpenCLEnvInternal(MagickCLEnv clEnv, ExceptionInfo* exception) { - MagickBooleanType status = MagickTrue; - cl_int clStatus; - cl_context_properties cps[3]; +static MagickCLDevice AcquireMagickCLDevice() +{ + MagickCLDevice + device; -#ifdef MAGICKCORE_CLPERFMARKER + device=(MagickCLDevice) AcquireMagickMemory(sizeof(*device)); + if (device != NULL) { - int status = clInitializePerfMarkerAMD(); - if (status == AP_SUCCESS) { - //printf("PerfMarker successfully initialized\n"); - } + (void) ResetMagickMemory(device,0,sizeof(*device)); + ActivateSemaphoreInfo(&device->lock); + device->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE; + device->command_queues_index=-1; + device->enabled=MagickTrue; } -#endif - clEnv->OpenCLInitialized = MagickTrue; - - /* check and init the global lib */ - OpenCLLib=GetOpenCLLib(); - if (OpenCLLib) - { - clEnv->library=OpenCLLib; - } - else - { - /* turn off opencl */ - MagickBooleanType flag; - flag = MagickTrue; - SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED - , sizeof(MagickBooleanType), &flag, exception); - } - - if (clEnv->OpenCLDisabled != MagickFalse) - goto cleanup; - - clEnv->OpenCLDisabled = MagickTrue; - /* setup the OpenCL platform and device */ - status = InitOpenCLPlatformDevice(clEnv, exception); - if (status == MagickFalse) { - /* No OpenCL device available */ - goto cleanup; - } - - /* create an OpenCL context */ - cps[0] = CL_CONTEXT_PLATFORM; - cps[1] = (cl_context_properties)clEnv->platform; - cps[2] = 0; - clEnv->context = clEnv->library->clCreateContext(cps, 1, &(clEnv->device), NULL, NULL, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "clCreateContext failed.", "(%d)", clStatus); - status = MagickFalse; - goto cleanup; - } - - status = CompileOpenCLKernels(clEnv, exception); - if (status == MagickFalse) { - (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning, - "clCreateCommandQueue failed.", "(%d)", status); - - status = MagickFalse; - goto cleanup; - } - - status = EnableOpenCLInternal(clEnv); - -cleanup: - return status; + return(device); } - -MagickExport -MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception) { - MagickBooleanType status = MagickFalse; - - if (clEnv == NULL) - return MagickFalse; - -#ifdef MAGICKCORE_CLPERFMARKER - clBeginPerfMarkerAMD(__FUNCTION__,""); -#endif - - LockSemaphoreInfo(clEnv->lock); - if (clEnv->OpenCLInitialized == MagickFalse) { - if (clEnv->device==NULL - && clEnv->OpenCLDisabled == MagickFalse) - status = autoSelectDevice(clEnv, exception); - else - status = InitOpenCLEnvInternal(clEnv, exception); - } - UnlockSemaphoreInfo(clEnv->lock); - -#ifdef MAGICKCORE_CLPERFMARKER - clEndPerfMarkerAMD(); -#endif - return status; -} - - /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % % % % % -+ A c q u i r e O p e n C L C o m m a n d Q u e u e % +% A c q u i r e M a g i c k C L E n v % +% % +% % +% % +%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% +% +% AcquireMagickCLEnv() allocates the MagickCLEnv structure +% +*/ + +static MagickCLEnv AcquireMagickCLEnv(void) +{ + const char + *option; + + MagickCLEnv + clEnv; + + clEnv=(MagickCLEnv) AcquireMagickMemory(sizeof(*clEnv)); + if (clEnv != (MagickCLEnv) NULL) + { + (void) ResetMagickMemory(clEnv,0,sizeof(*clEnv)); + ActivateSemaphoreInfo(&clEnv->lock); + clEnv->cpu_score=MAGICKCORE_OPENCL_UNDEFINED_SCORE; + clEnv->enabled=MagickTrue; + option=getenv("MAGICK_OCL_DEVICE"); + if ((option != (const char *) NULL) && (strcmp(option,"OFF") == 0)) + clEnv->enabled=MagickFalse; + } + return clEnv; +} + +/* +%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% +% % +% % +% % +% A c q u i r e O p e n C L C o m m a n d Q u e u e % % % % % % % @@ -1439,15 +549,15 @@ MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception) { % % The format of the AcquireOpenCLCommandQueue method is: % -% cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv) +% cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device) % % A description of each parameter follows: % -% o clEnv: the OpenCL environment. +% o device: the OpenCL device. % */ -MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv) +MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device) { cl_command_queue queue; @@ -1455,21 +565,23 @@ MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv) cl_command_queue_properties properties; - if (clEnv == (MagickCLEnv) NULL) - return (cl_command_queue) NULL; - LockSemaphoreInfo(clEnv->commandQueuesLock); - if (clEnv->commandQueuesPos >= 0) { - queue=clEnv->commandQueues[clEnv->commandQueuesPos--]; - UnlockSemaphoreInfo(clEnv->commandQueuesLock); + assert(device != (MagickCLDevice) NULL); + LockSemaphoreInfo(device->lock); + device->created_queues++; + if (device->command_queues_index >= 0) + { + queue=device->command_queues[device->command_queues_index--]; + UnlockSemaphoreInfo(device->lock); } - else { - UnlockSemaphoreInfo(clEnv->commandQueuesLock); - properties=0; -#if PROFILE_OCL_KERNELS + else + { + UnlockSemaphoreInfo(device->lock); + properties=(cl_command_queue_properties) NULL; +#if MAGICKCORE_OPENCL_PROFILE_KERNELS properties=CL_QUEUE_PROFILING_ENABLE; #endif - queue=clEnv->library->clCreateCommandQueue(clEnv->context,clEnv->device, - properties,NULL); + queue=openCL_library->clCreateCommandQueue(device->context, + device->deviceID,properties,NULL); } return(queue); } @@ -1479,59 +591,7 @@ MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv) % % % % % % -+ R e l i n q u i s h O p e n C L C o m m a n d Q u e u e % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% -% RelinquishOpenCLCommandQueue() releases the OpenCL command queue -% -% The format of the RelinquishOpenCLCommandQueue method is: -% -% MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv, -% cl_command_queue queue) -% -% A description of each parameter follows: -% -% o clEnv: the OpenCL environment. -% -% o queue: the OpenCL queue to be released. -% -% -*/ - -MagickPrivate MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv, - cl_command_queue queue) -{ - MagickBooleanType - status; - - if (clEnv == NULL) - return(MagickFalse); - - LockSemaphoreInfo(clEnv->commandQueuesLock); - - if (clEnv->commandQueuesPos >= MAX_COMMAND_QUEUES-1) - status=(clEnv->library->clReleaseCommandQueue(queue) == CL_SUCCESS) ? - MagickTrue : MagickFalse; - else - { - clEnv->commandQueues[++clEnv->commandQueuesPos]=queue; - status=MagickTrue; - } - - UnlockSemaphoreInfo(clEnv->commandQueuesLock); - - return(status); -} - -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -+ A c q u i r e O p e n C L K e r n e l % +% A c q u i r e O p e n C L K e r n e l % % % % % % % @@ -1554,62 +614,309 @@ MagickPrivate MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv, % */ -MagickPrivate cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv, - MagickOpenCLProgram program, const char* kernelName) +MagickPrivate cl_kernel AcquireOpenCLKernel(MagickCLDevice device, + const char *kernel_name) { cl_int - clStatus; + status; cl_kernel kernel; - kernel=(cl_kernel) NULL; - if ((clEnv != (MagickCLEnv) NULL) && - (kernelName != (MagickOpenCLProgram) NULL)) - { - kernel=clEnv->library->clCreateKernel(clEnv->programs[program],kernelName, - &clStatus); - } + assert(device != (MagickCLDevice) NULL); + kernel=openCL_library->clCreateKernel(device->program,kernel_name,&status); return(kernel); } - /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % % % % % -+ R e l i n q u i s h O p e n C L K e r n e l % +% A u t o S e l e c t O p e n C L D e v i c e s % % % % % % % %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % -% RelinquishOpenCLKernel() releases an OpenCL kernel +% AutoSelectOpenCLDevices() determines the best device based on the +% information from the micro-benchmark. % -% The format of the RelinquishOpenCLKernel method is: +% The format of the AutoSelectOpenCLDevices method is: % -% MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv, -% cl_kernel kernel) +% void AcquireOpenCLKernel(MagickCLEnv clEnv,ExceptionInfo *exception) % % A description of each parameter follows: % % o clEnv: the OpenCL environment. % -% o kernel: the OpenCL kernel object to be released. -% +% o exception: return any errors or warnings in this structure. % */ -MagickPrivate - MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv, cl_kernel kernel) +static void LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,const char *xml) { - MagickBooleanType status = MagickFalse; - if (clEnv != NULL && kernel != NULL) + char + keyword[MagickPathExtent], + *token; + + const char + *q; + + MagickCLDeviceBenchmark + *device_benchmark; + + MagickStatusType + status; + + size_t + i, + extent; + + if (xml == (char *) NULL) + return; + status=MagickTrue; + device_benchmark=(MagickCLDeviceBenchmark *) NULL; + token=AcquireString(xml); + extent=strlen(token)+MagickPathExtent; + for (q=(char *) xml; *q != '\0'; ) { - status = ((clEnv->library->clReleaseKernel(kernel) == CL_SUCCESS)?MagickTrue:MagickFalse); + /* + Interpret XML. + */ + GetNextToken(q,&q,extent,token); + if (*token == '\0') + break; + (void) CopyMagickString(keyword,token,MagickPathExtent); + if (LocaleNCompare(keyword,"",2) != 0) && (*q != '\0')) + GetNextToken(q,&q,extent,token); + continue; + } + if (LocaleNCompare(keyword,"