1/* 2 * Copyright (C) 2012 University of Szeged 3 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu> 4 * All rights reserved. 5 * 6 * Redistribution and use in source and binary forms, with or without 7 * modification, are permitted provided that the following conditions 8 * are met: 9 * 1. Redistributions of source code must retain the above copyright 10 * notice, this list of conditions and the following disclaimer. 11 * 2. Redistributions in binary form must reproduce the above copyright 12 * notice, this list of conditions and the following disclaimer in the 13 * documentation and/or other materials provided with the distribution. 14 * 15 * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY 16 * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE 17 * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR 18 * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL UNIVERSITY OF SZEGED OR 19 * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, 20 * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, 21 * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR 22 * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY 23 * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT 24 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE 25 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. 26 */ 27 28#include "config.h" 29 30#if ENABLE(OPENCL) 31#include "FilterContextOpenCL.h" 32 33namespace WebCore { 34 35FilterContextOpenCL* FilterContextOpenCL::m_context = 0; 36int FilterContextOpenCL::m_alreadyInitialized = 0; 37 38FilterContextOpenCL* FilterContextOpenCL::context() 39{ 40 if (m_context) 41 return m_context; 42 if (m_alreadyInitialized) 43 return 0; 44 45 m_alreadyInitialized = true; 46 FilterContextOpenCL* localContext = new FilterContextOpenCL(); 47 48 // Initializing the context. 49 cl_int errorNumber; 50 cl_device_id* devices; 51 cl_platform_id firstPlatformId; 52 size_t deviceBufferSize = 0; 53 54 errorNumber = clGetPlatformIDs(1, &firstPlatformId, 0); 55 cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)firstPlatformId, 0}; 56 localContext->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, 0, 0, &errorNumber); 57 if (errorNumber != CL_SUCCESS) { 58 localContext->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU, 0, 0, &errorNumber); 59 if (errorNumber != CL_SUCCESS) 60 return 0; 61 } 62 63 errorNumber = clGetContextInfo(localContext->m_deviceContext, CL_CONTEXT_DEVICES, 0, 0, &deviceBufferSize); 64 if (errorNumber != CL_SUCCESS) 65 return 0; 66 67 if (!deviceBufferSize) 68 return 0; 69 70 devices = reinterpret_cast<cl_device_id*>(fastMalloc(deviceBufferSize)); 71 errorNumber = clGetContextInfo(localContext->m_deviceContext, CL_CONTEXT_DEVICES, deviceBufferSize, devices, 0); 72 if (errorNumber != CL_SUCCESS) 73 return 0; 74 75 localContext->m_commandQueue = clCreateCommandQueue(localContext->m_deviceContext, devices[0], 0, 0); 76 if (!localContext->m_commandQueue) 77 return 0; 78 79 localContext->m_deviceId = devices[0]; 80 fastFree(devices); 81 82 cl_bool imageSupport = CL_FALSE; 83 clGetDeviceInfo(localContext->m_deviceId, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &imageSupport, 0); 84 if (imageSupport != CL_TRUE) 85 return 0; 86 87 m_context = localContext; 88 return m_context; 89} 90 91void FilterContextOpenCL::freeResources() 92{ 93 clFinish(m_commandQueue); 94 95 if (m_colorMatrixWasCompiled) { 96 freeResource(m_matrixOperation); 97 freeResource(m_saturateAndHueRotateOperation); 98 freeResource(m_luminanceOperation); 99 freeResource(m_colorMatrixProgram); 100 } 101 m_colorMatrixWasCompiled = false; 102 103 if (m_turbulenceWasCompiled) { 104 freeResource(m_turbulenceOperation); 105 freeResource(m_turbulenceProgram); 106 } 107 m_turbulenceWasCompiled = false; 108 109 if (m_transformColorSpaceWasCompiled) { 110 freeResource(m_transformColorSpaceKernel); 111 freeResource(m_transformColorSpaceProgram); 112 } 113 m_transformColorSpaceWasCompiled = false; 114} 115 116void FilterContextOpenCL::destroyContext() 117{ 118 freeResources(); 119 120 if (m_commandQueue) 121 clReleaseCommandQueue(m_commandQueue); 122 m_commandQueue = 0; 123 124 if (m_deviceContext) 125 clReleaseContext(m_deviceContext); 126 m_deviceContext = 0; 127 128 m_context = 0; 129} 130 131OpenCLHandle FilterContextOpenCL::createOpenCLImage(IntSize paintSize) 132{ 133 FilterContextOpenCL* context = FilterContextOpenCL::context(); 134 135 cl_image_format clImageFormat; 136 clImageFormat.image_channel_order = CL_RGBA; 137 clImageFormat.image_channel_data_type = CL_UNORM_INT8; 138 139#ifdef CL_API_SUFFIX__VERSION_1_2 140 cl_image_desc imageDescriptor = { CL_MEM_OBJECT_IMAGE2D, paintSize.width(), paintSize.height(), 0, 0, 0, 0, 0, 0, 0}; 141 OpenCLHandle image = clCreateImage(context->deviceContext(), CL_MEM_READ_WRITE, &clImageFormat, &imageDescriptor, 0, 0); 142#else 143 OpenCLHandle image = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE, &clImageFormat, 144 paintSize.width(), paintSize.height(), 0, 0, 0); 145#endif 146 return image; 147} 148 149static const char* transformColorSpaceKernelProgram = 150PROGRAM( 151const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; 152 153__kernel void transformColorSpace(__read_only image2d_t source, __write_only image2d_t destination, __constant float *clLookUpTable) 154{ 155 int2 sourceCoord = (int2) (get_global_id(0), get_global_id(1)); 156 float4 pixel = read_imagef(source, sampler, sourceCoord); 157 158 pixel = (float4) (clLookUpTable[(int)(round(pixel.x * 255))], clLookUpTable[(int)(round(pixel.y * 255))], 159 clLookUpTable[(int) (round(pixel.z * 255))], pixel.w); 160 161 write_imagef(destination, sourceCoord, pixel); 162} 163); // End of OpenCL kernels 164 165inline bool FilterContextOpenCL::compileTransformColorSpaceProgram() 166{ 167 if (m_transformColorSpaceWasCompiled || inError()) 168 return !inError(); 169 170 m_transformColorSpaceWasCompiled = true; 171 172 if (isResourceAllocationFailed((m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram)))) 173 return false; 174 if (isResourceAllocationFailed((m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace")))) 175 return false; 176 return true; 177} 178 179void FilterContextOpenCL::openCLTransformColorSpace(OpenCLHandle& source, IntRect sourceSize, ColorSpace srcColorSpace, ColorSpace dstColorSpace) 180{ 181 DEFINE_STATIC_LOCAL(OpenCLHandle, deviceRgbLUT, ()); 182 DEFINE_STATIC_LOCAL(OpenCLHandle, linearRgbLUT, ()); 183 184 if (srcColorSpace == dstColorSpace || inError()) 185 return; 186 187 if ((srcColorSpace != ColorSpaceLinearRGB && srcColorSpace != ColorSpaceDeviceRGB) 188 || (dstColorSpace != ColorSpaceLinearRGB && dstColorSpace != ColorSpaceDeviceRGB)) 189 return; 190 191 if (!compileTransformColorSpaceProgram()) 192 return; 193 194 OpenCLHandle destination = createOpenCLImage(sourceSize.size()); 195 196 RunKernel kernel(this, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height()); 197 kernel.addArgument(source); 198 kernel.addArgument(destination); 199 200 if (dstColorSpace == ColorSpaceLinearRGB) { 201 if (!linearRgbLUT) { 202 Vector<float> lookUpTable; 203 for (unsigned i = 0; i < 256; i++) { 204 float color = i / 255.0f; 205 color = (color <= 0.04045f ? color / 12.92f : pow((color + 0.055f) / 1.055f, 2.4f)); 206 color = std::max(0.0f, color); 207 color = std::min(1.0f, color); 208 lookUpTable.append((round(color * 255)) / 255); 209 } 210 linearRgbLUT = kernel.addArgument(lookUpTable.data(), sizeof(float) * 256); 211 } else 212 kernel.addArgument(linearRgbLUT); 213 } else if (dstColorSpace == ColorSpaceDeviceRGB) { 214 if (!deviceRgbLUT) { 215 Vector<float> lookUpTable; 216 for (unsigned i = 0; i < 256; i++) { 217 float color = i / 255.0f; 218 color = (powf(color, 1.0f / 2.4f) * 1.055f) - 0.055f; 219 color = std::max(0.0f, color); 220 color = std::min(1.0f, color); 221 lookUpTable.append((round(color * 255)) / 255); 222 } 223 deviceRgbLUT = kernel.addArgument(lookUpTable.data(), sizeof(float) * 256); 224 } else 225 kernel.addArgument(deviceRgbLUT); 226 } 227 228 kernel.run(); 229 source.clear(); 230 source = destination; 231} 232 233static const char* fillKernelProgram = 234PROGRAM_STR( 235__kernel void fill(__write_only image2d_t destination, float r, float g, float b, float a) 236{ 237 float4 sourcePixel = (float4)(r, g, b, a); 238 write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel); 239} 240); 241 242inline bool FilterContextOpenCL::compileFill() 243{ 244 if (m_fillWasCompiled || inError()) 245 return !inError(); 246 247 m_fillWasCompiled = true; 248 249 if (isResourceAllocationFailed((m_fillProgram = compileProgram(fillKernelProgram)))) 250 return false; 251 if (isResourceAllocationFailed((m_fill = kernelByName(m_fillProgram, "fill")))) 252 return false; 253 return true; 254} 255 256void FilterContextOpenCL::fill(cl_mem image, IntSize imageSize, Color color) 257{ 258 if (!m_context || inError()) 259 return; 260 261 compileFill(); 262 263 float r, g, b, a; 264 265 color.getRGBA(r, g, b, a); 266 267 RunKernel kernel(this, m_fill, imageSize.width(), imageSize.height()); 268 kernel.addArgument(image); 269 kernel.addArgument(r); 270 kernel.addArgument(g); 271 kernel.addArgument(b); 272 kernel.addArgument(a); 273 kernel.run(); 274} 275 276cl_program FilterContextOpenCL::compileProgram(const char* source) 277{ 278 cl_program program; 279 cl_int errorNumber = 0; 280 281 program = clCreateProgramWithSource(m_deviceContext, 1, (const char**) &source, 0, &errorNumber); 282 if (isFailed(errorNumber)) 283 return 0; 284 285 if (isFailed(clBuildProgram(program, 0, 0, 0, 0, 0))) 286 return 0; 287 288 return program; 289} 290 291void FilterContextOpenCL::freeResource(cl_kernel& handle) 292{ 293 if (handle) { 294 clReleaseKernel(handle); 295 handle = 0; 296 } 297} 298 299void FilterContextOpenCL::freeResource(cl_program& handle) 300{ 301 if (handle) { 302 clReleaseProgram(handle); 303 handle = 0; 304 } 305} 306} // namespace WebCore 307 308#endif 309