Changeset 142596 in webkit
- Timestamp:
- Feb 12, 2013 1:40:39 AM (11 years ago)
- Location:
- trunk/Source/WebCore
- Files:
-
- 8 edited
Legend:
- Unmodified
- Added
- Removed
-
trunk/Source/WebCore/ChangeLog
r142595 r142596 1 2013-02-12 Tamas Czene <tczene@inf.u-szeged.hu> 2 3 Add error checking into OpenCL version of SVG filters. 4 https://bugs.webkit.org/show_bug.cgi?id=107444 5 6 Reviewed by Zoltan Herczeg. 7 8 In case of an error the program runs through all the remaining filters by doing nothing. 9 After that deletes the results of every filter and starts software rendering. 10 11 * platform/graphics/filters/FilterEffect.cpp: 12 (WebCore): 13 (WebCore::FilterEffect::applyAll): At software rendering this is a simple inline methode, but at OpenCL rendering it releases OpenCL things. If we have an error remove filter's results and start software rendering. 14 (WebCore::FilterEffect::clearResultsRecursive): 15 (WebCore::FilterEffect::openCLImageToImageBuffer): 16 (WebCore::FilterEffect::createOpenCLImageResult): 17 (WebCore::FilterEffect::transformResultColorSpace): 18 * platform/graphics/filters/FilterEffect.h: 19 (FilterEffect): 20 (WebCore::FilterEffect::applyAll): 21 * platform/graphics/gpu/opencl/FilterContextOpenCL.cpp: 22 (WebCore::FilterContextOpenCL::isFailed): 23 (WebCore): 24 (WebCore::FilterContextOpenCL::freeResources): 25 (WebCore::FilterContextOpenCL::destroyContext): 26 (WebCore::FilterContextOpenCL::compileTransformColorSpaceProgram): 27 (WebCore::FilterContextOpenCL::openCLTransformColorSpace): 28 (WebCore::FilterContextOpenCL::compileProgram): 29 (WebCore::FilterContextOpenCL::freeResource): 30 * platform/graphics/gpu/opencl/FilterContextOpenCL.h: 31 (WebCore::FilterContextOpenCL::FilterContextOpenCL): 32 (WebCore::FilterContextOpenCL::setInError): 33 (WebCore::FilterContextOpenCL::inError): 34 (FilterContextOpenCL): 35 (WebCore::FilterContextOpenCL::RunKernel::RunKernel): 36 (WebCore::FilterContextOpenCL::RunKernel::addArgument): 37 (WebCore::FilterContextOpenCL::RunKernel::run): 38 (RunKernel): 39 * platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp: 40 (WebCore::FilterContextOpenCL::compileFEColorMatrix): 41 (WebCore::FEColorMatrix::platformApplyOpenCL): 42 * platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp: 43 (WebCore::FilterContextOpenCL::compileFETurbulence): 44 (WebCore::FETurbulence::platformApplyOpenCL): 45 * rendering/svg/RenderSVGResourceFilter.cpp: 46 (WebCore::RenderSVGResourceFilter::postApplyResource): 47 1 48 2013-02-12 Huang Dongsung <luxtella@company100.net> 2 49 -
trunk/Source/WebCore/platform/graphics/filters/FilterEffect.cpp
r142434 r142596 98 98 } 99 99 100 #if ENABLE(OPENCL) 101 void FilterEffect::applyAll() 102 { 103 if (hasResult()) 104 return; 105 FilterContextOpenCL* context = FilterContextOpenCL::context(); 106 if (context) { 107 apply(); 108 if (!context->inError()) 109 return; 110 clearResultsRecursive(); 111 context->destroyContext(); 112 } 113 // Software code path. 114 apply(); 115 } 116 #endif 117 100 118 void FilterEffect::apply() 101 119 { … … 224 242 } 225 243 244 void FilterEffect::clearResultsRecursive() 245 { 246 // Clear all results, regardless that the current effect has 247 // a result. Can be used if an effect is in an erroneous state. 248 if (hasResult()) 249 clearResult(); 250 251 unsigned size = m_inputEffects.size(); 252 for (unsigned i = 0; i < size; ++i) 253 m_inputEffects.at(i).get()->clearResultsRecursive(); 254 } 255 226 256 ImageBuffer* FilterEffect::asImageBuffer() 227 257 { … … 249 279 ASSERT(context); 250 280 281 if (context->inError()) 282 return 0; 283 251 284 size_t origin[3] = { 0, 0, 0 }; 252 285 size_t region[3] = { m_absolutePaintRect.width(), m_absolutePaintRect.height(), 1 }; … … 254 287 RefPtr<Uint8ClampedArray> destinationPixelArray = Uint8ClampedArray::create(m_absolutePaintRect.width() * m_absolutePaintRect.height() * 4); 255 288 256 clFinish(context->commandQueue()); 257 clEnqueueReadImage(context->commandQueue(), m_openCLImageResult, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0); 289 if (context->isFailed(clFinish(context->commandQueue()))) 290 return 0; 291 292 if (context->isFailed(clEnqueueReadImage(context->commandQueue(), m_openCLImageResult, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0))) 293 return 0; 258 294 259 295 m_imageBufferResult = ImageBuffer::create(m_absolutePaintRect.size()); … … 427 463 OpenCLHandle FilterEffect::createOpenCLImageResult(uint8_t* source) 428 464 { 465 FilterContextOpenCL* context = FilterContextOpenCL::context(); 466 ASSERT(context); 467 468 if (context->inError()) 469 return 0; 470 429 471 ASSERT(!hasResult()); 430 472 cl_image_format clImageFormat; … … 432 474 clImageFormat.image_channel_data_type = CL_UNORM_INT8; 433 475 434 FilterContextOpenCL* context = FilterContextOpenCL::context(); 435 ASSERT(context); 476 int errorCode = 0; 436 477 m_openCLImageResult = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE | (source ? CL_MEM_COPY_HOST_PTR : 0), 437 &clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, 0); 478 &clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, &errorCode); 479 if (context->isFailed(errorCode)) 480 return 0; 481 438 482 return m_openCLImageResult; 439 483 } … … 451 495 #if ENABLE(OPENCL) 452 496 if (openCLImage()) { 497 if (m_imageBufferResult) 498 m_imageBufferResult.clear(); 453 499 FilterContextOpenCL* context = FilterContextOpenCL::context(); 454 500 ASSERT(context); 455 501 context->openCLTransformColorSpace(m_openCLImageResult, absolutePaintRect(), m_resultColorSpace, dstColorSpace); 456 if (m_imageBufferResult)457 m_imageBufferResult.clear();458 502 } else { 459 503 #endif 460 504 461 // FIXME: We can avoid this potentially unnecessary ImageBuffer conversion by adding462 // color space transform support for the {pre,un}multiplied arrays.463 asImageBuffer()->transformColorSpace(m_resultColorSpace, dstColorSpace);505 // FIXME: We can avoid this potentially unnecessary ImageBuffer conversion by adding 506 // color space transform support for the {pre,un}multiplied arrays. 507 asImageBuffer()->transformColorSpace(m_resultColorSpace, dstColorSpace); 464 508 465 509 #if ENABLE(OPENCL) -
trunk/Source/WebCore/platform/graphics/filters/FilterEffect.h
r138250 r142596 69 69 70 70 void clearResult(); 71 void clearResultsRecursive(); 72 71 73 ImageBuffer* asImageBuffer(); 72 74 PassRefPtr<Uint8ClampedArray> asUnmultipliedImage(const IntRect&); … … 110 112 111 113 void apply(); 112 114 #if ENABLE(OPENCL) 115 void applyAll(); 116 #else 117 inline void applyAll() { apply(); } 118 #endif 119 113 120 // Correct any invalid pixels, if necessary, in the result of a filter operation. 114 121 // This method is used to ensure valid pixel values on filter inputs and the final result. -
trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp
r137999 r142596 89 89 } 90 90 91 void 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 116 void 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 91 131 OpenCLHandle FilterContextOpenCL::createOpenCLImage(IntSize paintSize) 92 132 { … … 118 158 ); // End of OpenCL kernels 119 159 160 inline bool FilterContextOpenCL::compileTransformColorSpaceProgram() 161 { 162 if (m_transformColorSpaceWasCompiled || inError()) 163 return !inError(); 164 165 m_transformColorSpaceWasCompiled = true; 166 167 if (isResourceAllocationFailed((m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram)))) 168 return false; 169 if (isResourceAllocationFailed((m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace")))) 170 return false; 171 return true; 172 } 173 120 174 void FilterContextOpenCL::openCLTransformColorSpace(OpenCLHandle& source, IntRect sourceSize, ColorSpace srcColorSpace, ColorSpace dstColorSpace) 121 175 { … … 123 177 DEFINE_STATIC_LOCAL(OpenCLHandle, linearRgbLUT, ()); 124 178 125 if (srcColorSpace == dstColorSpace )179 if (srcColorSpace == dstColorSpace || inError()) 126 180 return; 127 181 … … 130 184 return; 131 185 132 FilterContextOpenCL* context = FilterContextOpenCL::context(); 133 ASSERT(context); 134 135 OpenCLHandle destination = context->createOpenCLImage(sourceSize.size()); 136 137 if (!m_transformColorSpaceProgram) { 138 m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram); 139 ASSERT(m_transformColorSpaceProgram); 140 m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace"); 141 ASSERT(m_transformColorSpaceKernel); 142 } 143 144 RunKernel kernel(context, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height()); 186 if (!compileTransformColorSpaceProgram()) 187 return; 188 189 OpenCLHandle destination = createOpenCLImage(sourceSize.size()); 190 191 RunKernel kernel(this, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height()); 145 192 kernel.addArgument(source); 146 193 kernel.addArgument(destination); … … 182 229 { 183 230 cl_program program; 184 cl_int errorNumber; 185 186 FilterContextOpenCL* context = FilterContextOpenCL::context(); 187 ASSERT(context); 188 189 program = clCreateProgramWithSource(context->m_deviceContext, 1, (const char**) &source, 0, 0); 190 errorNumber = clBuildProgram(program, 0, 0, 0, 0, 0); 191 if (errorNumber) 231 cl_int errorNumber = 0; 232 233 program = clCreateProgramWithSource(m_deviceContext, 1, (const char**) &source, 0, &errorNumber); 234 if (isFailed(errorNumber)) 235 return 0; 236 237 if (isFailed(clBuildProgram(program, 0, 0, 0, 0, 0))) 192 238 return 0; 193 239 194 240 return program; 195 241 } 242 243 void FilterContextOpenCL::freeResource(cl_kernel& handle) 244 { 245 if (handle) { 246 clReleaseKernel(handle); 247 handle = 0; 248 } 249 } 250 251 void FilterContextOpenCL::freeResource(cl_program& handle) 252 { 253 if (handle) { 254 clReleaseProgram(handle); 255 handle = 0; 256 } 257 } 196 258 } // namespace WebCore 197 259 -
trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h
r137999 r142596 45 45 public: 46 46 FilterContextOpenCL() 47 : m_deviceId(0) 47 : m_inError(false) 48 , m_deviceId(0) 48 49 , m_deviceContext(0) 49 50 , m_commandQueue(0) 51 , m_transformColorSpaceWasCompiled(false) 50 52 , m_transformColorSpaceProgram(0) 51 53 , m_transformColorSpaceKernel(0) 52 , m_colorMatrix CompileStatus(openclNotCompiledYet)54 , m_colorMatrixWasCompiled(false) 53 55 , m_colorMatrixProgram(0) 54 56 , m_matrixOperation(0) 55 57 , m_saturateAndHueRotateOperation(0) 56 58 , m_luminanceOperation(0) 57 , m_turbulence CompileStatus(openclNotCompiledYet)59 , m_turbulenceWasCompiled(false) 58 60 , m_turbulenceProgram(0) 59 61 , m_turbulenceOperation(0) … … 68 70 cl_command_queue commandQueue() { return m_commandQueue; } 69 71 72 inline void setInError(bool errorCode = true) { m_inError = errorCode; } 73 inline bool inError() { return m_inError; } 74 inline bool isFailed(bool); 75 inline bool isResourceAllocationFailed(bool); 76 77 void freeResources(); 78 void destroyContext(); 79 70 80 OpenCLHandle createOpenCLImage(IntSize); 81 82 inline bool compileTransformColorSpaceProgram(); 71 83 void openCLTransformColorSpace(OpenCLHandle&, IntRect, ColorSpace, ColorSpace); 72 84 … … 85 97 : m_context(context) 86 98 , m_kernel(kernel) 87 , index(0) 99 , m_index(0) 100 , m_error(context->inError()) 88 101 { 89 102 m_globalSize[0] = width; … … 93 106 void addArgument(OpenCLHandle handle) 94 107 { 95 clSetKernelArg(m_kernel, index++, sizeof(OpenCLHandle), handle.handleAddress()); 108 if (!m_error) 109 m_error = clSetKernelArg(m_kernel, m_index++, sizeof(OpenCLHandle), handle.handleAddress()); 96 110 } 97 111 98 112 void addArgument(cl_int value) 99 113 { 100 clSetKernelArg(m_kernel, index++, sizeof(cl_int), reinterpret_cast<void*>(&value)); 114 if (!m_error) 115 m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_int), reinterpret_cast<void*>(&value)); 101 116 } 102 117 103 118 void addArgument(cl_float value) 104 119 { 105 clSetKernelArg(m_kernel, index++, sizeof(cl_float), reinterpret_cast<void*>(&value)); 120 if (!m_error) 121 m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_float), reinterpret_cast<void*>(&value)); 106 122 } 107 123 108 124 void addArgument(cl_sampler handle) 109 125 { 110 clSetKernelArg(m_kernel, index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle)); 126 if (!m_error) 127 m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle)); 111 128 } 112 129 113 130 OpenCLHandle addArgument(void* buffer, int size) 114 131 { 115 OpenCLHandle handle(clCreateBuffer(m_context->deviceContext(), CL_MEM_READ_ONLY, size, 0, 0)); 116 clEnqueueWriteBuffer(m_context->commandQueue(), handle, CL_TRUE, 0, size, buffer, 0, 0, 0); 117 clSetKernelArg(m_kernel, index++, sizeof(OpenCLHandle), handle.handleAddress()); 118 return handle; 132 if (m_error) 133 return 0; 134 OpenCLHandle handle(clCreateBuffer(m_context->deviceContext(), CL_MEM_READ_ONLY, size, 0, &m_error)); 135 if (m_error) 136 return 0; 137 m_error = clEnqueueWriteBuffer(m_context->commandQueue(), handle, CL_TRUE, 0, size, buffer, 0, 0, 0); 138 if (m_error) 139 return 0; 140 m_error = clSetKernelArg(m_kernel, m_index++, sizeof(OpenCLHandle), handle.handleAddress()); 141 return !m_error ? handle : 0; 119 142 } 120 143 121 144 void run() 122 145 { 123 clFinish(m_context->m_commandQueue); 124 clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0); 146 if (m_context->isFailed(m_error)) 147 return; 148 149 m_error = clFinish(m_context->m_commandQueue); 150 if (!m_error) 151 m_error = clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0); 152 m_context->isFailed(m_error); 125 153 } 126 154 … … 128 156 cl_kernel m_kernel; 129 157 size_t m_globalSize[2]; 130 int index; 158 int m_index; 159 int m_error; 131 160 }; 132 161 133 enum OpenCLCompileStatus { 134 openclNotCompiledYet, 135 openclCompileFailed, 136 openclCompileSuccessful 137 }; 138 139 static cl_program compileProgram(const char*); 162 cl_program compileProgram(const char*); 140 163 static inline cl_kernel kernelByName(cl_program program, const char* name) { return clCreateKernel(program, name, 0); } 164 165 static inline void freeResource(cl_kernel&); 166 static inline void freeResource(cl_program&); 141 167 142 168 static FilterContextOpenCL* m_context; 143 169 static int m_alreadyInitialized; 170 bool m_inError; 144 171 145 172 cl_device_id m_deviceId; … … 147 174 cl_command_queue m_commandQueue; 148 175 176 bool m_transformColorSpaceWasCompiled; 149 177 cl_program m_transformColorSpaceProgram; 150 178 cl_kernel m_transformColorSpaceKernel; 151 179 152 OpenCLCompileStatus m_colorMatrixCompileStatus;180 bool m_colorMatrixWasCompiled; 153 181 cl_program m_colorMatrixProgram; 154 182 cl_kernel m_matrixOperation; … … 156 184 cl_kernel m_luminanceOperation; 157 185 158 OpenCLCompileStatus m_turbulenceCompileStatus;186 bool m_turbulenceWasCompiled; 159 187 cl_program m_turbulenceProgram; 160 188 cl_kernel m_turbulenceOperation; 161 189 }; 162 190 191 inline bool FilterContextOpenCL::isFailed(bool value) 192 { 193 if (value) 194 setInError(); 195 return value; 196 } 197 198 inline bool FilterContextOpenCL::isResourceAllocationFailed(bool value) 199 { 200 if (!value) 201 setInError(); 202 return !value; 203 } 204 163 205 } // namespace WebCore 164 206 -
trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp
r137999 r142596 69 69 inline bool FilterContextOpenCL::compileFEColorMatrix() 70 70 { 71 if (m_colorMatrix CompileStatus != openclNotCompiledYet)72 return m_colorMatrixCompileStatus == openclCompileSuccessful;71 if (m_colorMatrixWasCompiled || inError()) 72 return !inError(); 73 73 74 m_colorMatrix CompileStatus = openclCompileFailed;75 m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram); 76 if ( !m_colorMatrixProgram)74 m_colorMatrixWasCompiled = true; 75 76 if (isResourceAllocationFailed((m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram)))) 77 77 return false; 78 79 m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix"); 80 if (!m_matrixOperation) 78 if (isResourceAllocationFailed((m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix")))) 81 79 return false; 82 m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate"); 83 if (!m_saturateAndHueRotateOperation) 80 if (isResourceAllocationFailed((m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate")))) 84 81 return false; 85 m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance"); 86 if (!m_luminanceOperation) 82 if (isResourceAllocationFailed((m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate")))) 87 83 return false; 88 89 m_colorMatrixCompileStatus = openclCompileSuccessful;90 return openclCompileSuccessful;84 if (isResourceAllocationFailed((m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance")))) 85 return false; 86 return true; 91 87 } 92 88 … … 131 127 { 132 128 FilterContextOpenCL* context = FilterContextOpenCL::context(); 133 if (!context || !context->compileFEColorMatrix())129 if (!context) 134 130 return false; 131 132 if (!context->compileFEColorMatrix()) 133 return true; 135 134 136 135 FilterEffect* in = inputEffect(0); -
trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp
r137999 r142596 170 170 inline bool FilterContextOpenCL::compileFETurbulence() 171 171 { 172 if (m_turbulence CompileStatus != openclNotCompiledYet)173 return m_turbulenceCompileStatus == openclCompileSuccessful;174 175 m_turbulence CompileStatus = openclCompileFailed;176 m_turbulenceProgram = compileProgram(turbulenceKernelProgram); 177 if ( !m_turbulenceProgram)172 if (m_turbulenceWasCompiled || inError()) 173 return !inError(); 174 175 m_turbulenceWasCompiled = true; 176 177 if (isResourceAllocationFailed((m_turbulenceProgram = compileProgram(turbulenceKernelProgram)))) 178 178 return false; 179 m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence"); 180 if (!m_turbulenceOperation) 179 if (isResourceAllocationFailed((m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence")))) 181 180 return false; 182 183 m_turbulenceCompileStatus = openclCompileSuccessful; 184 return openclCompileSuccessful; 181 return true; 185 182 } 186 183 … … 225 222 { 226 223 FilterContextOpenCL* context = FilterContextOpenCL::context(); 227 if (!context || !context->compileFETurbulence())224 if (!context) 228 225 return false; 226 227 if (!context->compileFETurbulence()) 228 return true; 229 229 230 230 OpenCLHandle destination = createOpenCLImageResult(); -
trunk/Source/WebCore/rendering/svg/RenderSVGResourceFilter.cpp
r138835 r142596 310 310 if (!lastEffect->hasResult()) { 311 311 filterData->state = FilterData::Applying; 312 lastEffect->apply ();312 lastEffect->applyAll(); 313 313 lastEffect->correctFilterResultIfNeeded(); 314 314 lastEffect->transformResultColorSpace(ColorSpaceDeviceRGB);
Note: See TracChangeset
for help on using the changeset viewer.