From 66acef5cd2089f66bbdb7dc7b3b18e2eb6d792ae Mon Sep 17 00:00:00 2001 From: dirk Date: Tue, 21 Jun 2016 00:02:08 +0200 Subject: [PATCH] Added queue parameter to EnqueueOpenCLKernel. The command queue is now flushed when it is being released. Minor refactoring. --- MagickCore/accelerate.c | 178 ++++++++++++++++++++++++------------ MagickCore/cache.c | 3 +- MagickCore/opencl-private.h | 24 ++--- MagickCore/opencl.c | 152 ++++++++++++++---------------- 4 files changed, 204 insertions(+), 153 deletions(-) diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index 85aed3511..3a341c44e 100644 --- a/MagickCore/accelerate.c +++ b/MagickCore/accelerate.c @@ -371,6 +371,9 @@ cleanup: static Image *ComputeAddNoiseImage(const Image *image,MagickCLEnv clEnv, const NoiseType noise_type,ExceptionInfo *exception) { + cl_command_queue + queue; + cl_float attenuate; @@ -423,6 +426,9 @@ static Image *ComputeAddNoiseImage(const Image *image,MagickCLEnv clEnv, outputReady=MagickFalse; device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); + if (queue == (cl_command_queue) NULL) + goto cleanup; imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -510,12 +516,15 @@ static Image *ComputeAddNoiseImage(const Image *image,MagickCLEnv clEnv, goto cleanup; } - outputReady=EnqueueOpenCLKernel(addNoiseKernel,1,(const size_t *) NULL,gsize, + outputReady=EnqueueOpenCLKernel(queue,addNoiseKernel,1,(const size_t *) NULL,gsize, lsize,image,filteredImage,exception); + cleanup: if (addNoiseKernel != (cl_kernel) NULL) ReleaseOpenCLKernel(addNoiseKernel); + if (queue != (cl_command_queue) NULL) + ReleaseOpenCLCommandQueue(device,queue); if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL)) @@ -562,6 +571,9 @@ MagickPrivate Image *AccelerateAddNoiseImage(const Image *image, static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv, const double radius,const double sigma,ExceptionInfo *exception) { + cl_command_queue + queue; + cl_int status; @@ -607,6 +619,7 @@ static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv, outputReady=MagickFalse; device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -625,7 +638,7 @@ static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv, length=image->columns*image->rows; tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length* - sizeof(cl_float4),NULL); + sizeof(cl_float4),(void *) NULL); if (tempImageBuffer == (cl_mem) NULL) goto cleanup; @@ -636,13 +649,6 @@ static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv, ResourceLimitWarning,"AcquireOpenCLKernel failed.","."); goto cleanup; } - blurColumnKernel=AcquireOpenCLKernel(device,"BlurColumn"); - if (blurColumnKernel == (cl_kernel) NULL) - { - (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), - ResourceLimitWarning,"AcquireOpenCLKernel failed.","."); - goto cleanup; - } number_channels=(cl_uint) image->number_channels; imageColumns=(cl_uint) image->columns; @@ -670,11 +676,19 @@ static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv, lsize[0]=chunkSize; lsize[1]=1; - outputReady=EnqueueOpenCLKernel(blurRowKernel,2,NULL,gsize,lsize,image, - filteredImage,exception); + outputReady=EnqueueOpenCLKernel(queue,blurRowKernel,2,(size_t *) NULL,gsize, + lsize,image,filteredImage,exception); if (outputReady == MagickFalse) goto cleanup; + blurColumnKernel=AcquireOpenCLKernel(device,"BlurColumn"); + if (blurColumnKernel == (cl_kernel) NULL) + { + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), + ResourceLimitWarning,"AcquireOpenCLKernel failed.","."); + goto cleanup; + } + i=0; status =SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); status|=SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_uint),&number_channels); @@ -697,8 +711,8 @@ static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv, lsize[0]=1; lsize[1]=chunkSize; - outputReady=EnqueueOpenCLKernel(blurColumnKernel,2,NULL,gsize,lsize,image, - filteredImage,exception); + outputReady=EnqueueOpenCLKernel(queue,blurColumnKernel,2,(size_t *) NULL,gsize, + lsize,image,filteredImage,exception); cleanup: @@ -710,6 +724,8 @@ cleanup: ReleaseOpenCLKernel(blurRowKernel); if (blurColumnKernel != (cl_kernel) NULL) ReleaseOpenCLKernel(blurColumnKernel); + if (queue != (cl_command_queue) NULL) + ReleaseOpenCLCommandQueue(device,queue); if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL)) @@ -889,7 +905,7 @@ cleanup: if (filterKernel!=NULL) ReleaseOpenCLKernel(filterKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(device,queue); + ReleaseOpenCLCommandQueue(device,queue); if (device != NULL) ReleaseOpenCLDevice(device); @@ -1489,7 +1505,7 @@ cleanup: if (stretchKernel!=NULL) ReleaseOpenCLKernel(stretchKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(device,queue); + ReleaseOpenCLCommandQueue(device,queue); if (device != NULL) ReleaseOpenCLDevice(device); @@ -1843,7 +1859,7 @@ cleanup: if (clkernel != NULL) ReleaseOpenCLKernel(clkernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(device,queue); + ReleaseOpenCLCommandQueue(device,queue); if (device != NULL) ReleaseOpenCLDevice(device); if (outputReady == MagickFalse) @@ -2236,7 +2252,7 @@ cleanup: filteredImage_view=DestroyCacheView(filteredImage_view); if (queue != NULL) - RelinquishOpenCLCommandQueue(device,queue); + ReleaseOpenCLCommandQueue(device,queue); if (device != NULL) ReleaseOpenCLDevice(device); if (imageBuffer!=NULL) @@ -2712,7 +2728,7 @@ cleanup: if (equalizeKernel!=NULL) ReleaseOpenCLKernel(equalizeKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(device, queue); + ReleaseOpenCLCommandQueue(device, queue); if (device != NULL) ReleaseOpenCLDevice(device); @@ -2759,6 +2775,9 @@ static MagickBooleanType ComputeFunctionImage(Image *image,MagickCLEnv clEnv, const MagickFunction function,const size_t number_parameters, const double *parameters,ExceptionInfo *exception) { + cl_command_queue + queue; + cl_int status; @@ -2792,6 +2811,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image,MagickCLEnv clEnv, parametersBuffer=NULL; device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -2840,7 +2860,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image,MagickCLEnv clEnv, gsize[0]=image->columns; gsize[1]=image->rows; - outputReady=EnqueueOpenCLKernel(functionKernel,2,(const size_t *) NULL, + outputReady=EnqueueOpenCLKernel(queue,functionKernel,2,(const size_t *) NULL, gsize,(const size_t *) NULL,image,(const Image *) NULL,exception); cleanup: @@ -2849,6 +2869,8 @@ cleanup: ReleaseOpenCLMemObject(parametersBuffer); if (functionKernel != (cl_kernel) NULL) ReleaseOpenCLKernel(functionKernel); + if (queue != (cl_command_queue) NULL) + ReleaseOpenCLCommandQueue(device,queue); if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); return(outputReady); @@ -2894,6 +2916,9 @@ MagickPrivate MagickBooleanType AccelerateFunctionImage(Image *image, static MagickBooleanType ComputeGrayscaleImage(Image *image,MagickCLEnv clEnv, const PixelIntensityMethod method,ExceptionInfo *exception) { + cl_command_queue + queue; + cl_int status; @@ -2924,6 +2949,7 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image,MagickCLEnv clEnv, assert(image != (Image *) NULL); assert(image->signature == MagickCoreSignature); device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -2954,13 +2980,16 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image,MagickCLEnv clEnv, gsize[0]=image->columns; gsize[1]=image->rows; - outputReady=EnqueueOpenCLKernel(grayscaleKernel,2,(const size_t *) NULL, - gsize,(const size_t *) NULL,image,(Image *) NULL,exception); + outputReady=EnqueueOpenCLKernel(queue,grayscaleKernel,2, + (const size_t *) NULL,gsize,(const size_t *) NULL,image,(Image *) NULL, + exception); cleanup: if (grayscaleKernel != (cl_kernel) NULL) ReleaseOpenCLKernel(grayscaleKernel); + if (queue != (cl_command_queue) NULL) + ReleaseOpenCLCommandQueue(device,queue); if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); @@ -3309,7 +3338,7 @@ cleanup: if (blurColumnKernel!=NULL) ReleaseOpenCLKernel(blurColumnKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(device, queue); + ReleaseOpenCLCommandQueue(device, queue); if (device != NULL) ReleaseOpenCLDevice(device); if (outputReady == MagickFalse) @@ -3522,7 +3551,7 @@ cleanup: if (modulateKernel!=NULL) ReleaseOpenCLKernel(modulateKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(device,queue); + ReleaseOpenCLCommandQueue(device,queue); if (device != NULL) ReleaseOpenCLDevice(device); @@ -3890,7 +3919,7 @@ cleanup: if (motionBlurKernel!=NULL) ReleaseOpenCLKernel(motionBlurKernel); if (queue != NULL) - RelinquishOpenCLCommandQueue(device,queue); + ReleaseOpenCLCommandQueue(device,queue); if (device != NULL) ReleaseOpenCLDevice(device); if (outputReady == MagickFalse && filteredImage != NULL) @@ -3939,8 +3968,8 @@ MagickPrivate Image *AccelerateMotionBlurImage(const Image *image, */ static MagickBooleanType resizeHorizontalFilter(MagickCLDevice device, - const Image *image,Image *filteredImage,cl_mem imageBuffer, - cl_uint number_channels,cl_uint columns,cl_uint rows, + cl_command_queue queue,const Image *image,Image *filteredImage, + cl_mem imageBuffer,cl_uint number_channels,cl_uint columns,cl_uint rows, cl_mem resizedImageBuffer,cl_uint resizedColumns,cl_uint resizedRows, const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients, const float xFactor,ExceptionInfo *exception) @@ -4110,8 +4139,8 @@ RestoreMSCWarning gsize[1]=resizedRows; lsize[0]=workgroupSize; lsize[1]=1; - outputReady=EnqueueOpenCLKernel(horizontalKernel,2,(const size_t *) NULL, - gsize,lsize,image,filteredImage,exception); + outputReady=EnqueueOpenCLKernel(queue,horizontalKernel,2, + (const size_t *) NULL,gsize,lsize,image,filteredImage,exception); cleanup: if (horizontalKernel != (cl_kernel) NULL) @@ -4121,8 +4150,8 @@ cleanup: } static MagickBooleanType resizeVerticalFilter(MagickCLDevice device, - const Image *image,Image * filteredImage,cl_mem imageBuffer, - cl_uint number_channels,cl_uint columns,cl_uint rows, + cl_command_queue queue,const Image *image,Image * filteredImage, + cl_mem imageBuffer,cl_uint number_channels,cl_uint columns,cl_uint rows, cl_mem resizedImageBuffer,cl_uint resizedColumns,cl_uint resizedRows, const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients, const float yFactor,ExceptionInfo *exception) @@ -4292,7 +4321,7 @@ RestoreMSCWarning workgroupSize; lsize[0]=1; lsize[1]=workgroupSize; - outputReady=EnqueueOpenCLKernel(verticalKernel,2,(const size_t *) NULL, + outputReady=EnqueueOpenCLKernel(queue,verticalKernel,2,(const size_t *) NULL, gsize,lsize,image,filteredImage,exception); cleanup: @@ -4307,6 +4336,9 @@ static Image *ComputeResizeImage(const Image* image,MagickCLEnv clEnv, const size_t resizedColumns,const size_t resizedRows, const ResizeFilter *resizeFilter,ExceptionInfo *exception) { + cl_command_queue + queue; + cl_mem cubicCoefficientsBuffer, filteredImageBuffer, @@ -4345,6 +4377,7 @@ static Image *ComputeResizeImage(const Image* image,MagickCLEnv clEnv, outputReady=MagickFalse; device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -4359,8 +4392,8 @@ static Image *ComputeResizeImage(const Image* image,MagickCLEnv clEnv, resizeFilterCoefficient=GetResizeFilterCoefficient(resizeFilter); for (i = 0; i < 7; i++) coefficientBuffer[i]=(float) resizeFilterCoefficient[i]; - cubicCoefficientsBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY | - CL_MEM_COPY_HOST_PTR,7*sizeof(float),&coefficientBuffer); + cubicCoefficientsBuffer=CreateOpenCLBuffer(device,CL_MEM_COPY_HOST_PTR | + CL_MEM_READ_ONLY,7*sizeof(*resizeFilterCoefficient),&coefficientBuffer); if (cubicCoefficientsBuffer == (cl_mem) NULL) { (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), @@ -4383,17 +4416,19 @@ static Image *ComputeResizeImage(const Image* image,MagickCLEnv clEnv, goto cleanup; } - outputReady=resizeHorizontalFilter(device,image,filteredImage,imageBuffer, - number_channels,(cl_uint) image->columns,(cl_uint) image->rows, - tempImageBuffer,(cl_uint) resizedColumns,(cl_uint) image->rows, - resizeFilter,cubicCoefficientsBuffer,xFactor,exception); + outputReady=resizeHorizontalFilter(device,queue,image,filteredImage, + imageBuffer,number_channels,(cl_uint) image->columns, + (cl_uint) image->rows,tempImageBuffer,(cl_uint) resizedColumns, + (cl_uint) image->rows,resizeFilter,cubicCoefficientsBuffer,xFactor, + exception); if (outputReady == MagickFalse) goto cleanup; - outputReady=resizeVerticalFilter(device,image,filteredImage,tempImageBuffer, - number_channels,(cl_uint) resizedColumns,(cl_uint) image->rows, - filteredImageBuffer,(cl_uint) resizedColumns,(cl_uint) resizedRows, - resizeFilter,cubicCoefficientsBuffer,yFactor,exception); + outputReady=resizeVerticalFilter(device,queue,image,filteredImage, + tempImageBuffer,number_channels,(cl_uint) resizedColumns, + (cl_uint) image->rows,filteredImageBuffer,(cl_uint) resizedColumns, + (cl_uint) resizedRows,resizeFilter,cubicCoefficientsBuffer,yFactor, + exception); if (outputReady == MagickFalse) goto cleanup; } @@ -4409,17 +4444,19 @@ static Image *ComputeResizeImage(const Image* image,MagickCLEnv clEnv, goto cleanup; } - outputReady=resizeVerticalFilter(device,image,filteredImage,imageBuffer, - number_channels,(cl_uint) image->columns,(cl_int) image->rows, - tempImageBuffer,(cl_uint) image->columns,(cl_uint) resizedRows, - resizeFilter,cubicCoefficientsBuffer,yFactor,exception); + outputReady=resizeVerticalFilter(device,queue,image,filteredImage, + imageBuffer,number_channels,(cl_uint) image->columns, + (cl_int) image->rows,tempImageBuffer,(cl_uint) image->columns, + (cl_uint) resizedRows,resizeFilter,cubicCoefficientsBuffer,yFactor, + exception); if (outputReady == MagickFalse) goto cleanup; - outputReady=resizeHorizontalFilter(device,image,filteredImage,tempImageBuffer, - number_channels,(cl_uint) image->columns, (cl_uint) resizedRows, - filteredImageBuffer,(cl_uint) resizedColumns, (cl_uint) resizedRows, - resizeFilter,cubicCoefficientsBuffer,xFactor,exception); + outputReady=resizeHorizontalFilter(device,queue,image,filteredImage, + tempImageBuffer,number_channels,(cl_uint) image->columns, + (cl_uint) resizedRows,filteredImageBuffer,(cl_uint) resizedColumns, + (cl_uint) resizedRows,resizeFilter,cubicCoefficientsBuffer,xFactor, + exception); if (outputReady == MagickFalse) goto cleanup; } @@ -4430,6 +4467,8 @@ cleanup: ReleaseOpenCLMemObject(tempImageBuffer); if (cubicCoefficientsBuffer != (cl_mem) NULL) ReleaseOpenCLMemObject(cubicCoefficientsBuffer); + if (queue != (cl_command_queue) NULL) + ReleaseOpenCLCommandQueue(device,queue); if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL)) @@ -4500,6 +4539,9 @@ MagickPrivate Image *AccelerateResizeImage(const Image *image, static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv, const double angle,ExceptionInfo *exception) { + cl_command_queue + queue; + cl_float2 blurCenter; @@ -4552,6 +4594,7 @@ static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv, outputReady=MagickFalse; device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -4634,8 +4677,9 @@ static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv, gsize[0]=image->columns; gsize[1]=image->rows; - outputReady=EnqueueOpenCLKernel(rotationalBlurKernel,2,(const size_t *) NULL, - gsize,(const size_t *) NULL,image,filteredImage,exception); + outputReady=EnqueueOpenCLKernel(queue,rotationalBlurKernel,2, + (const size_t *) NULL,gsize,(const size_t *) NULL,image,filteredImage, + exception); cleanup: @@ -4645,6 +4689,8 @@ cleanup: ReleaseOpenCLMemObject(cosThetaBuffer); if (rotationalBlurKernel != (cl_kernel) NULL) ReleaseOpenCLKernel(rotationalBlurKernel); + if (queue != (cl_command_queue) NULL) + ReleaseOpenCLCommandQueue(device,queue); if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL)) @@ -4692,6 +4738,9 @@ static Image *ComputeUnsharpMaskImage(const Image *image,MagickCLEnv clEnv, const double radius,const double sigma,const double gain, const double threshold,ExceptionInfo *exception) { + cl_command_queue + queue; + cl_int status; @@ -4743,6 +4792,7 @@ static Image *ComputeUnsharpMaskImage(const Image *image,MagickCLEnv clEnv, outputReady=MagickFalse; device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -4811,8 +4861,8 @@ static Image *ComputeUnsharpMaskImage(const Image *image,MagickCLEnv clEnv, gsize[1]=image->rows; lsize[0]=chunkSize; lsize[1]=1; - outputReady=EnqueueOpenCLKernel(blurRowKernel,2,(const size_t *) NULL,gsize, - lsize,image,filteredImage,exception); + outputReady=EnqueueOpenCLKernel(queue,blurRowKernel,2, + (const size_t *) NULL,gsize,lsize,image,filteredImage,exception); chunkSize=256; fGain=(float) gain; @@ -4843,7 +4893,7 @@ static Image *ComputeUnsharpMaskImage(const Image *image,MagickCLEnv clEnv, gsize[1]=chunkSize*((image->rows+chunkSize-1)/chunkSize); lsize[0]=1; lsize[1]=chunkSize; - outputReady=EnqueueOpenCLKernel(unsharpMaskBlurColumnKernel,2, + outputReady=EnqueueOpenCLKernel(queue,unsharpMaskBlurColumnKernel,2, (const size_t *) NULL,gsize,lsize,image,filteredImage,exception); cleanup: @@ -4856,6 +4906,8 @@ cleanup: ReleaseOpenCLKernel(blurRowKernel); if (unsharpMaskBlurColumnKernel != (cl_kernel) NULL) ReleaseOpenCLKernel(unsharpMaskBlurColumnKernel); + if (queue != (cl_command_queue) NULL) + ReleaseOpenCLCommandQueue(device,queue); if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL)) @@ -4868,6 +4920,9 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, MagickCLEnv clEnv,const double radius,const double sigma,const double gain, const double threshold,ExceptionInfo *exception) { + cl_command_queue + queue; + cl_int status; @@ -4909,6 +4964,7 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, outputReady=MagickFalse; device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -4960,7 +5016,7 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image, gsize[1]=((image->rows + 31) / 32)*32; lsize[0]=8; lsize[1]=32; - outputReady=EnqueueOpenCLKernel(unsharpMaskKernel,2,(const size_t *) NULL, + outputReady=EnqueueOpenCLKernel(queue,unsharpMaskKernel,2,(const size_t *) NULL, gsize,lsize,image,filteredImage,exception); cleanup: @@ -4969,6 +5025,8 @@ cleanup: ReleaseOpenCLMemObject(imageKernelBuffer); if (unsharpMaskKernel != (cl_kernel) NULL) ReleaseOpenCLKernel(unsharpMaskKernel); + if (queue != (cl_command_queue) NULL) + ReleaseOpenCLCommandQueue(device,queue); if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL)) @@ -5009,6 +5067,9 @@ MagickPrivate Image *AccelerateUnsharpMaskImage(const Image *image, static Image *ComputeWaveletDenoiseImage(const Image *image,MagickCLEnv clEnv, const double threshold,ExceptionInfo *exception) { + cl_command_queue + queue; + const cl_int PASSES=5; @@ -5055,6 +5116,7 @@ static Image *ComputeWaveletDenoiseImage(const Image *image,MagickCLEnv clEnv, outputReady=MagickFalse; device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); if (imageBuffer == (cl_mem) NULL) goto cleanup; @@ -5102,13 +5164,15 @@ static Image *ComputeWaveletDenoiseImage(const Image *image,MagickCLEnv clEnv, gsize[1]=((height+(SIZE-1))/SIZE)*4; lsize[0]=TILESIZE; lsize[1]=4; - outputReady=EnqueueOpenCLKernel(denoiseKernel,2,(const size_t *) NULL,gsize, - lsize,image,filteredImage,exception); + outputReady=EnqueueOpenCLKernel(queue,denoiseKernel,2,(const size_t *) NULL, + gsize,lsize,image,filteredImage,exception); cleanup: if (denoiseKernel != (cl_kernel) NULL) ReleaseOpenCLKernel(denoiseKernel); + if (queue != (cl_command_queue) NULL) + ReleaseOpenCLCommandQueue(device,queue); if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL)) diff --git a/MagickCore/cache.c b/MagickCore/cache.c index 0793a3333..1d4e6596c 100644 --- a/MagickCore/cache.c +++ b/MagickCore/cache.c @@ -1167,12 +1167,13 @@ MagickPrivate cl_mem GetAuthenticOpenCLBuffer(const Image *image, cache_info->opencl=CopyMagickCLCacheInfo(cache_info->opencl); if (cache_info->opencl == (MagickCLCacheInfo) NULL) { - assert(cache_info->pixels != NULL); + assert(cache_info->pixels != (Quantum *) NULL); cache_info->opencl=AcquireMagickCLCacheInfo(device,cache_info->pixels, cache_info->length); if (cache_info->opencl == (MagickCLCacheInfo) NULL) return((cl_mem) NULL); } + assert(cache_info->opencl->pixels == cache_info->pixels); return(cache_info->opencl->buffer); } #endif diff --git a/MagickCore/opencl-private.h b/MagickCore/opencl-private.h index 49190b2a1..fd61e195b 100644 --- a/MagickCore/opencl-private.h +++ b/MagickCore/opencl-private.h @@ -105,6 +105,14 @@ typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseCommandQueue)( cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0; +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; + /* Memory Object APIs */ typedef CL_API_ENTRY cl_mem @@ -222,12 +230,6 @@ typedef CL_API_ENTRY cl_int 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; - -/* Finish APIs */ -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 @@ -243,6 +245,8 @@ struct MagickLibraryRec MAGICKpfn_clCreateCommandQueue clCreateCommandQueue; MAGICKpfn_clReleaseCommandQueue clReleaseCommandQueue; + MAGICKpfn_clFlush clFlush; + MAGICKpfn_clFinish clFinish; MAGICKpfn_clCreateBuffer clCreateBuffer; MAGICKpfn_clReleaseMemObject clReleaseMemObject; @@ -270,8 +274,6 @@ struct MagickLibraryRec MAGICKpfn_clSetEventCallback clSetEventCallback; MAGICKpfn_clGetEventProfilingInfo clGetEventProfilingInfo; - - MAGICKpfn_clFinish clFinish; }; struct _MagickCLDevice @@ -407,8 +409,8 @@ extern MagickPrivate cl_mem CreateOpenCLBuffer(MagickCLDevice,cl_mem_flags,size_t,void *); extern MagickPrivate MagickBooleanType - EnqueueOpenCLKernel(cl_kernel,cl_uint,const size_t *,const size_t *, - const size_t *,const Image *,const Image *,ExceptionInfo *), + EnqueueOpenCLKernel(cl_command_queue,cl_kernel,cl_uint,const size_t *, + const size_t *,const size_t *,const Image *,const Image *,ExceptionInfo *), InitializeOpenCL(MagickCLEnv,ExceptionInfo *), OpenCLThrowMagickException(MagickCLDevice,ExceptionInfo *, const char *,const char *,const size_t,const ExceptionType,const char *, @@ -432,10 +434,10 @@ extern MagickPrivate void DumpOpenCLProfileData(), OpenCLTerminus(), RecordProfileData(MagickCLDevice,cl_kernel,cl_event), + ReleaseOpenCLCommandQueue(MagickCLDevice,cl_command_queue), ReleaseOpenCLDevice(MagickCLDevice), ReleaseOpenCLKernel(cl_kernel), ReleaseOpenCLMemObject(cl_mem), - RelinquishOpenCLCommandQueue(MagickCLDevice,cl_command_queue), RetainOpenCLEvent(cl_event); #endif diff --git a/MagickCore/opencl.c b/MagickCore/opencl.c index a740f19cd..8d8908ce5 100644 --- a/MagickCore/opencl.c +++ b/MagickCore/opencl.c @@ -438,7 +438,7 @@ static size_t StringSignature(const char* string) */ MagickPrivate cl_mem CreateOpenCLBuffer(MagickCLDevice device, - cl_mem_flags flags, size_t size, void *host_ptr) + cl_mem_flags flags,size_t size,void *host_ptr) { return(openCL_library->clCreateBuffer(device->context,flags,size,host_ptr, (cl_int *) NULL)); @@ -508,7 +508,8 @@ MagickPrivate MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device, info->length=length; info->pixels=pixels; info->buffer=openCL_library->clCreateBuffer(device->context, - CL_MEM_USE_HOST_PTR,(size_t) length,(void *) pixels,&status); + CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,(size_t) length,(void *) pixels, + &status); if (status == CL_SUCCESS) return(info); LockSemaphoreInfo(openCL_lock); @@ -1400,9 +1401,6 @@ MagickPrivate MagickCLCacheInfo CopyMagickCLCacheInfo(MagickCLCacheInfo info) cl_command_queue queue; - cl_event - event; - Quantum *pixels; @@ -1411,13 +1409,11 @@ MagickPrivate MagickCLCacheInfo CopyMagickCLCacheInfo(MagickCLCacheInfo info) if (info->event_count > 0) { queue=AcquireOpenCLCommandQueue(info->device); - pixels=openCL_library->clEnqueueMapBuffer(queue,info->buffer,CL_FALSE, + pixels=openCL_library->clEnqueueMapBuffer(queue,info->buffer,CL_TRUE, CL_MAP_READ | CL_MAP_WRITE,0,info->length,info->event_count, - info->events,&event,(cl_int *) NULL); + info->events,(cl_event *) NULL,(cl_int *) NULL); assert(pixels == info->pixels); - RelinquishOpenCLCommandQueue(info->device,queue); - openCL_library->clWaitForEvents(1,&event); - openCL_library->clReleaseEvent(event); + ReleaseOpenCLCommandQueue(info->device,queue); } return(RelinquishMagickCLCacheInfo(info,MagickFalse)); } @@ -1578,8 +1574,8 @@ static void RegisterCacheEvent(MagickCLCacheInfo info,cl_event event) openCL_library->clRetainEvent(event); } -MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel, - cl_uint work_dim,const size_t *offset,const size_t *gsize, +MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue, + cl_kernel kernel,cl_uint work_dim,const size_t *offset,const size_t *gsize, const size_t *lsize,const Image *input_image,const Image *output_image, ExceptionInfo *exception) { @@ -1587,9 +1583,6 @@ MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel, *output_info, *input_info; - cl_command_queue - queue; - cl_event event, *events; @@ -1604,9 +1597,6 @@ MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel, input_info=(CacheInfo *) input_image->cache; assert(input_info != (CacheInfo *) NULL); assert(input_info->opencl != (MagickCLCacheInfo) NULL); - queue=AcquireOpenCLCommandQueue(input_info->opencl->device); - if (queue == (cl_command_queue) NULL) - return(MagickFalse); event_count=input_info->opencl->event_count; events=input_info->opencl->events; output_info=(CacheInfo *) NULL; @@ -1623,10 +1613,7 @@ MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel, event_count+=output_info->opencl->event_count; events=AcquireQuantumMemory(event_count,sizeof(*events)); if (events == (cl_event *) NULL) - { - RelinquishOpenCLCommandQueue(input_info->opencl->device,queue); - return(MagickFalse); - } + return(MagickFalse); for (i=0; i < (ssize_t) event_count; i++) { if (i < (ssize_t) input_info->opencl->event_count) @@ -1639,7 +1626,6 @@ MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel, } status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset, gsize,lsize,event_count,events,&event); - RelinquishOpenCLCommandQueue(input_info->opencl->device,queue); if ((output_info != (CacheInfo *) NULL) && (output_info->opencl->event_count > 0)) events=(cl_event *) RelinquishMagickMemory(events); @@ -2364,6 +2350,8 @@ static MagickBooleanType BindOpenCLFunctions() BIND(clCreateCommandQueue); BIND(clReleaseCommandQueue); + BIND(clFlush); + BIND(clFinish); BIND(clCreateProgramWithSource); BIND(clCreateProgramWithBinary); @@ -2389,8 +2377,6 @@ static MagickBooleanType BindOpenCLFunctions() BIND(clGetEventProfilingInfo); - BIND(clFinish); - return(MagickTrue); } @@ -2580,13 +2566,13 @@ MagickPrivate void RecordProfileData(MagickCLDevice device, if (status != CL_SUCCESS) return; name=AcquireQuantumMemory(length,sizeof(*name)); - (void) openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length, - name,NULL); + status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length, + name,(size_t *) NULL); start=end=elapsed=0; - openCL_library->clWaitForEvents(1,&event); - status=openCL_library->clGetEventProfilingInfo(event, + status|=openCL_library->clWaitForEvents(1,&event); + status|=openCL_library->clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL); - status&=openCL_library->clGetEventProfilingInfo(event, + status|=openCL_library->clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL); if (status != CL_SUCCESS) { @@ -2611,13 +2597,13 @@ MagickPrivate void RecordProfileData(MagickCLDevice device, i++; } } - if (profile_record == ((KernelProfileRecord) NULL)) + if (profile_record == (KernelProfileRecord) NULL) { profile_record=AcquireMagickMemory(sizeof(*profile_record)); (void) ResetMagickMemory(profile_record,0,sizeof(*profile_record)); profile_record->kernel_name=AcquireString(name); device->profile_records=ResizeMagickMemory(device->profile_records,(i+2)* - sizeof(KernelProfileRecord)); + sizeof(*device->profile_records)); device->profile_records[i]=profile_record; device->profile_records[i+1]=(KernelProfileRecord) NULL; } @@ -2631,6 +2617,54 @@ MagickPrivate void RecordProfileData(MagickCLDevice device, name=DestroyString(name); } +/* +%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% +% % +% % +% % ++ R e l e a s e O p e n C L C o m m a n d Q u e u e % +% % +% % +% % +%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% +% +% ReleaseOpenCLCommandQueue() releases the OpenCL command queue +% +% The format of the ReleaseOpenCLCommandQueue method is: +% +% void ReleaseOpenCLCommandQueue(MagickCLDevice device, +% cl_command_queue queue) +% +% A description of each parameter follows: +% +% o device: the OpenCL device. +% +% o queue: the OpenCL queue to be released. +*/ + +MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device, + cl_command_queue queue) +{ + if (queue == (cl_command_queue) NULL) + return; + + assert(device != (MagickCLDevice) NULL); + LockSemaphoreInfo(device->lock); + if ((device->profile_kernels != MagickFalse) || + (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1)) + { + UnlockSemaphoreInfo(device->lock); + openCL_library->clFinish(queue); + (void) openCL_library->clReleaseCommandQueue(queue); + } + else + { + openCL_library->clFlush(queue); + device->command_queues[++device->command_queues_index]=queue; + UnlockSemaphoreInfo(device->lock); + } +} + /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % @@ -2695,12 +2729,9 @@ static void DestroyMagickCLCacheInfo(MagickCLCacheInfo info) for (i=0; i < (ssize_t) info->event_count; i++) openCL_library->clReleaseEvent(info->events[i]); - info->events=RelinquishMagickMemory(info->events); + info->events=(cl_event *) RelinquishMagickMemory(info->events); if (info->buffer != (cl_mem) NULL) - { - openCL_library->clReleaseMemObject(info->buffer); - info->buffer=(cl_mem) NULL; - } + openCL_library->clReleaseMemObject(info->buffer); ReleaseOpenCLDevice(info->device); RelinquishMagickMemory(info); } @@ -2820,53 +2851,6 @@ static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv) return((MagickCLEnv) RelinquishMagickMemory(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: -% -% void RelinquishOpenCLCommandQueue(MagickCLDevice device, -% cl_command_queue queue) -% -% A description of each parameter follows: -% -% o device: the OpenCL device. -% -% o queue: the OpenCL queue to be released. -*/ - -MagickPrivate void RelinquishOpenCLCommandQueue(MagickCLDevice device, - cl_command_queue queue) -{ - if (queue == (cl_command_queue) NULL) - return; - - assert(device != (MagickCLDevice) NULL); - LockSemaphoreInfo(device->lock); - if ((device->profile_kernels != MagickFalse) || - (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1)) - { - UnlockSemaphoreInfo(device->lock); - openCL_library->clFinish(queue); - (void) openCL_library->clReleaseCommandQueue(queue); - } - else - { - device->command_queues[++device->command_queues_index]=queue; - UnlockSemaphoreInfo(device->lock); - } -} - /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % -- 2.40.0