return(MagickFalse);
/* check if the image has mask */
- if (image->composite_mask != MagickFalse ||
- image->read_mask != MagickFalse ||
- image->write_mask != MagickFalse)
+ if (((image->channels & ReadMaskChannel) != 0) ||
+ ((image->channels & WriteMaskChannel) != 0) ||
+ ((image->channels & CompositeMaskChannel) != 0))
return(MagickFalse);
if (image->number_channels > 4)
return(clone);
}
-/* pad the global workgroup size to the next multiple of
+/* pad the global workgroup size to the next multiple of
the local workgroup size */
inline static unsigned int padGlobalWorkgroupSizeToLocalWorkgroupSize(
- const unsigned int orgGlobalSize,const unsigned int localGroupSize)
+ const unsigned int orgGlobalSize,const unsigned int localGroupSize)
{
return ((orgGlobalSize+(localGroupSize-1))/localGroupSize*localGroupSize);
}
size_t
global_work_size[2];
- histogramKernel = NULL;
+ histogramKernel = NULL;
outputReady = MagickFalse;
colorspace = image->colorspace;
outputReady = MagickTrue;
cleanup:
-
+
if (histogramKernel!=NULL)
ReleaseOpenCLKernel(histogramKernel);
imageBuffer = NULL;
histogramBuffer = NULL;
stretchMapBuffer = NULL;
- histogramKernel = NULL;
- stretchKernel = NULL;
+ histogramKernel = NULL;
+ stretchKernel = NULL;
queue = NULL;
outputReady = MagickFalse;
if (histogram == (cl_uint4 *) NULL)
ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed", image->filename);
-
+
/* reset histogram */
(void) memset(histogram,0,(MaxMap+1)*sizeof(*histogram));
(void) OpenCLThrowMagickException(device,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,
+ /* 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))
+ if (ALIGNED(inputPixels,CLPixelPacket))
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
}
- else
+ else
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
}
goto cleanup;
}
- /* If the host pointer is aligned to the size of cl_uint,
- then use the host buffer directly from the GPU; otherwise,
+ /* If the host pointer is aligned to the size of cl_uint,
+ then use the host buffer directly from the GPU; otherwise,
create a buffer on the GPU and copy the data over */
- if (ALIGNED(histogram,cl_uint4))
+ if (ALIGNED(histogram,cl_uint4))
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
hostPtr = histogram;
}
- else
+ else
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
hostPtr = histogram;
}
/* create a CL buffer for histogram */
- length = (MaxMap+1);
+ length = (MaxMap+1);
histogramBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus);
if (clStatus != CL_SUCCESS)
{
goto cleanup;
/* read from the kenel output */
- if (ALIGNED(histogram,cl_uint4))
+ if (ALIGNED(histogram,cl_uint4))
{
- length = (MaxMap+1);
+ length = (MaxMap+1);
clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
}
- else
+ else
{
- length = (MaxMap+1);
+ length = (MaxMap+1);
clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length * sizeof(cl_uint4), histogram, 0, NULL, NULL);
}
if (clStatus != CL_SUCCESS)
}
/* recreate input buffer later, in case image updated */
-#ifdef RECREATEBUFFER
- if (imageBuffer!=NULL)
+#ifdef RECREATEBUFFER
+ if (imageBuffer!=NULL)
clEnv->library->clReleaseMemObject(imageBuffer);
#endif
if (stretch_map == (PixelPacket *) NULL)
ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed",
image->filename);
-
+
/*
Stretch the histogram to create the stretched image mapping.
*/
stretch_map: uchar4 (PixelPacket)
black, white: float4 (FloatPixelPacket) */
-#ifdef RECREATEBUFFER
- /* If the host pointer is aligned to the size of CLPixelPacket,
- then use the host buffer directly from the GPU; otherwise,
+#ifdef RECREATEBUFFER
+ /* 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))
+ if (ALIGNED(inputPixels,CLPixelPacket))
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
}
- else
+ else
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
}
#endif
/* Create and initialize OpenCL buffers. */
- if (ALIGNED(stretch_map, PixelPacket))
+ if (ALIGNED(stretch_map, PixelPacket))
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
hostPtr = stretch_map;
}
- else
+ else
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
hostPtr = stretch_map;
}
/* create a CL buffer for stretch_map */
- length = (MaxMap+1);
+ length = (MaxMap+1);
stretchMapBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus);
if (clStatus != CL_SUCCESS)
{
RecordProfileData(device,stretchKernel,event);
/* read the data back */
- if (ALIGNED(inputPixels,CLPixelPacket))
+ 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
+ else
{
length = image->columns * image->rows;
clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
image_view=DestroyCacheView(image_view);
- if (imageBuffer!=NULL)
+ if (imageBuffer!=NULL)
clEnv->library->clReleaseMemObject(imageBuffer);
if (stretchMapBuffer!=NULL)
/* Create and initialize OpenCL buffers. */
- /* If the host pointer is aligned to the size of CLPixelPacket,
- then use the host buffer directly from the GPU; otherwise,
+ /* 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))
+ if (ALIGNED(inputPixels,CLPixelPacket))
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
}
- else
+ else
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
}
goto cleanup;
}
- if (ALIGNED(filteredPixels,CLPixelPacket))
+ if (ALIGNED(filteredPixels,CLPixelPacket))
{
mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
hostPtr = filteredPixels;
}
- else
+ else
{
mem_flags = CL_MEM_WRITE_ONLY;
hostPtr = NULL;
global_work_size[0] = (image->columns + (localGroupSize[0]-1))/localGroupSize[0] * localGroupSize[0];
global_work_size[1] = (image->rows + (localGroupSize[1]-1))/localGroupSize[1] * localGroupSize[1];
clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, &event);
-
+
if (clStatus != CL_SUCCESS)
{
(void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
}
RecordProfileData(device,clkernel,event);
- if (ALIGNED(filteredPixels,CLPixelPacket))
+ if (ALIGNED(filteredPixels,CLPixelPacket))
{
length = image->columns * image->rows;
clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
}
- else
+ else
{
length = image->columns * image->rows;
clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
ExceptionInfo*exception)
{
- static const int
+ static const int
X[4] = {0, 1, 1,-1},
Y[4] = {1, 0, 1, 1};
device = RequestOpenCLDevice(clEnv);
queue = AcquireOpenCLCommandQueue(device);
-
+
image_view=AcquireAuthenticCacheView(image,exception);
inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
if (inputPixels == (void *) NULL)
goto cleanup;
}
- if (ALIGNED(inputPixels,CLPixelPacket))
+ if (ALIGNED(inputPixels,CLPixelPacket))
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
}
- else
+ else
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
}
goto cleanup;
}
- if (ALIGNED(filteredPixels,CLPixelPacket))
+ if (ALIGNED(filteredPixels,CLPixelPacket))
{
mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
hostPtr = filteredPixels;
}
- else
+ else
{
mem_flags = CL_MEM_WRITE_ONLY;
hostPtr = NULL;
global_work_size[0] = image->columns;
global_work_size[1] = image->rows;
-
+
for (k = 0; k < 4; k++)
{
cl_int2 offset;
int polarity;
-
+
offset.s[0] = X[k];
offset.s[1] = Y[k];
polarity = 1;
RecordProfileData(device,hullPass2,event);
}
- if (ALIGNED(filteredPixels,CLPixelPacket))
+ if (ALIGNED(filteredPixels,CLPixelPacket))
{
length = image->columns * image->rows;
clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
}
- else
+ else
{
length = image->columns * image->rows;
clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
imageBuffer = NULL;
histogramBuffer = NULL;
equalizeMapBuffer = NULL;
- histogramKernel = NULL;
- equalizeKernel = NULL;
+ histogramKernel = NULL;
+ equalizeKernel = NULL;
queue = NULL;
outputReady = MagickFalse;
(void) OpenCLThrowMagickException(device,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,
+ /* 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))
+ if (ALIGNED(inputPixels,CLPixelPacket))
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
}
- else
+ else
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
}
goto cleanup;
}
- /* If the host pointer is aligned to the size of cl_uint,
- then use the host buffer directly from the GPU; otherwise,
+ /* If the host pointer is aligned to the size of cl_uint,
+ then use the host buffer directly from the GPU; otherwise,
create a buffer on the GPU and copy the data over */
- if (ALIGNED(histogram,cl_uint4))
+ if (ALIGNED(histogram,cl_uint4))
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
hostPtr = histogram;
}
- else
+ else
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
hostPtr = histogram;
}
/* create a CL buffer for histogram */
- length = (MaxMap+1);
+ length = (MaxMap+1);
histogramBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus);
if (clStatus != CL_SUCCESS)
{
goto cleanup;
/* read from the kenel output */
- if (ALIGNED(histogram,cl_uint4))
+ if (ALIGNED(histogram,cl_uint4))
{
- length = (MaxMap+1);
+ length = (MaxMap+1);
clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
}
- else
+ else
{
- length = (MaxMap+1);
+ length = (MaxMap+1);
clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length * sizeof(cl_uint4), histogram, 0, NULL, NULL);
}
if (clStatus != CL_SUCCESS)
}
/* recreate input buffer later, in case image updated */
-#ifdef RECREATEBUFFER
- if (imageBuffer!=NULL)
+#ifdef RECREATEBUFFER
+ if (imageBuffer!=NULL)
clEnv->library->clReleaseMemObject(imageBuffer);
#endif
-
+
/* CPU stuff */
equalize_map=(PixelPacket *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*equalize_map));
if (equalize_map == (PixelPacket *) NULL)
equalize_map: uchar4 (PixelPacket)
black, white: float4 (FloatPixelPacket) */
-#ifdef RECREATEBUFFER
- /* If the host pointer is aligned to the size of CLPixelPacket,
- then use the host buffer directly from the GPU; otherwise,
+#ifdef RECREATEBUFFER
+ /* 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))
+ if (ALIGNED(inputPixels,CLPixelPacket))
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
}
- else
+ else
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
}
#endif
/* Create and initialize OpenCL buffers. */
- if (ALIGNED(equalize_map, PixelPacket))
+ if (ALIGNED(equalize_map, PixelPacket))
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
hostPtr = equalize_map;
}
- else
+ else
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
hostPtr = equalize_map;
}
/* create a CL buffer for eqaulize_map */
- length = (MaxMap+1);
+ length = (MaxMap+1);
equalizeMapBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus);
if (clStatus != CL_SUCCESS)
{
RecordProfileData(device,equalizeKernel,event);
/* read the data back */
- if (ALIGNED(inputPixels,CLPixelPacket))
+ 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
+ else
{
length = image->columns * image->rows;
clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
goto cleanup;
}
- /* If the host pointer is aligned to the size of CLPixelPacket,
- then use the host buffer directly from the GPU; otherwise,
+ /* 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))
+ if (ALIGNED(inputPixels,CLPixelPacket))
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
}
- else
+ else
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
}
goto cleanup;
}
- if (ALIGNED(filteredPixels,CLPixelPacket))
+ if (ALIGNED(filteredPixels,CLPixelPacket))
{
mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
hostPtr = filteredPixels;
}
- else
+ else
{
mem_flags = CL_MEM_WRITE_ONLY;
hostPtr = NULL;
clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_int),(void *)&iRadius);
clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
-
+
if (clStatus != CL_SUCCESS)
{
(void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
}
/* get result */
- if (ALIGNED(filteredPixels,CLPixelPacket))
+ if (ALIGNED(filteredPixels,CLPixelPacket))
{
length = image->columns * image->rows;
clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
}
- else
+ else
{
length = image->columns * image->rows;
clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
inputPixels = NULL;
imageBuffer = NULL;
- modulateKernel = NULL;
+ modulateKernel = NULL;
assert(image != (Image *) NULL);
assert(image->signature == MagickCoreSignature);
goto cleanup;
}
- /* If the host pointer is aligned to the size of CLPixelPacket,
- then use the host buffer directly from the GPU; otherwise,
+ /* 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))
+ if (ALIGNED(inputPixels,CLPixelPacket))
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
}
- else
+ else
{
mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
}
RecordProfileData(device,modulateKernel,event);
}
- if (ALIGNED(inputPixels,CLPixelPacket))
+ 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
+ else
{
length = image->columns * image->rows;
clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
cl_mem
filteredImageBuffer,
imageBuffer,
- imageKernelBuffer,
+ imageKernelBuffer,
offsetBuffer;
cl_mem_flags
goto cleanup;
}
- // If the host pointer is aligned to the size of CLPixelPacket,
- // then use the host buffer directly from the GPU; otherwise,
+ // 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))
+ if (ALIGNED(inputPixels,CLPixelPacket))
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
}
- else
+ else
{
mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
}
// create a CL buffer from image pixel buffer
length = image->columns * image->rows;
- imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags,
+ imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags,
length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
if (clStatus != CL_SUCCESS)
{
assert(filteredImage != NULL);
if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
{
- (void) ThrowMagickException(exception, GetMagickModule(),
+ (void) ThrowMagickException(exception, GetMagickModule(),
ResourceLimitError, "CloneImage failed.", ".");
goto cleanup;
}
filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
if (filteredPixels == (void *) NULL)
{
- (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
+ (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
"UnableToReadPixelCache.","`%s'",filteredImage->filename);
goto cleanup;
}
- if (ALIGNED(filteredPixels,CLPixelPacket))
+ if (ALIGNED(filteredPixels,CLPixelPacket))
{
mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
hostPtr = filteredPixels;
}
- else
+ else
{
mem_flags = CL_MEM_WRITE_ONLY;
hostPtr = NULL;
}
// create a CL buffer from image pixel buffer
length = image->columns * image->rows;
- filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags,
+ filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags,
length * sizeof(CLPixelPacket), hostPtr, &clStatus);
if (clStatus != CL_SUCCESS)
{
- (void) ThrowMagickException(exception, GetMagickModule(),
+ (void) ThrowMagickException(exception, GetMagickModule(),
ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
goto cleanup;
}
- imageKernelBuffer = clEnv->library->clCreateBuffer(device->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)
{
- (void) ThrowMagickException(exception, GetMagickModule(),
+ (void) ThrowMagickException(exception, GetMagickModule(),
ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
goto cleanup;
}
queue = AcquireOpenCLCommandQueue(device);
- kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer,
+ kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer,
CL_TRUE, CL_MAP_WRITE, 0, width * sizeof(float), 0, NULL, NULL, &clStatus);
if (clStatus != CL_SUCCESS)
{
- (void) ThrowMagickException(exception, GetMagickModule(),
+ (void) ThrowMagickException(exception, GetMagickModule(),
ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.",".");
goto cleanup;
}
0, NULL, NULL);
if (clStatus != CL_SUCCESS)
{
- (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
+ (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
"clEnv->library->clEnqueueUnmapMemObject failed.", ".");
goto cleanup;
}
- offsetBuffer = clEnv->library->clCreateBuffer(device->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)
{
- (void) ThrowMagickException(exception, GetMagickModule(),
+ (void) ThrowMagickException(exception, GetMagickModule(),
ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
goto cleanup;
}
- offsetBufferPtr = (int*)clEnv->library->clEnqueueMapBuffer(queue, offsetBuffer, CL_TRUE,
+ offsetBufferPtr = (int*)clEnv->library->clEnqueueMapBuffer(queue, offsetBuffer, CL_TRUE,
CL_MAP_WRITE, 0, width * sizeof(cl_int2), 0, NULL, NULL, &clStatus);
if (clStatus != CL_SUCCESS)
{
- (void) ThrowMagickException(exception, GetMagickModule(),
+ (void) ThrowMagickException(exception, GetMagickModule(),
ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.",".");
goto cleanup;
}
offsetBufferPtr[2*i] = (int)offset[i].x;
offsetBufferPtr[2*i+1] = (int)offset[i].y;
}
- clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, offsetBuffer, offsetBufferPtr, 0,
+ clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, offsetBuffer, offsetBufferPtr, 0,
NULL, NULL);
if (clStatus != CL_SUCCESS)
{
"AcquireOpenCLKernel failed.", ".");
goto cleanup;
}
-
+
// set the kernel arguments
i = 0;
clStatus=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
(unsigned int) image->columns,(unsigned int) local_work_size[0]);
global_work_size[1] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
(unsigned int) image->rows,(unsigned int) local_work_size[1]);
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, motionBlurKernel, 2, NULL,
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, motionBlurKernel, 2, NULL,
global_work_size, local_work_size, 0, NULL, &event);
if (clStatus != CL_SUCCESS)
}
RecordProfileData(device,motionBlurKernel,event);
- if (ALIGNED(filteredPixels,CLPixelPacket))
+ if (ALIGNED(filteredPixels,CLPixelPacket))
{
length = image->columns * image->rows;
- clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE,
- CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL,
+ clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE,
+ CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL,
NULL, &clStatus);
}
- else
+ else
{
length = image->columns * image->rows;
- clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0,
+ clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0,
length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
}
if (clStatus != CL_SUCCESS)
}
scale=PerceptibleReciprocal(scale);
- if (resizedColumns < workgroupSize)
+ if (resizedColumns < workgroupSize)
{
chunkSize=32;
pixelPerWorkgroup=32;
}
scale=PerceptibleReciprocal(scale);
- if (resizedRows < workgroupSize)
+ if (resizedRows < workgroupSize)
{
chunkSize=32;
pixelPerWorkgroup=32;
MagickPrivate Image *AccelerateResizeImage(const Image *image,
const size_t resizedColumns,const size_t resizedRows,
- const ResizeFilter *resizeFilter,ExceptionInfo *exception)
+ const ResizeFilter *resizeFilter,ExceptionInfo *exception)
{
Image
*filteredImage;