https://bugs.webkit.org/show_bug.cgi?id=107444
Patch by Tamas Czene <tczene@inf.u-szeged.hu> on 2013-02-12
Reviewed by Zoltan Herczeg.
In case of an error the program runs through all the remaining filters by doing nothing.
After that deletes the results of every filter and starts software rendering.
* platform/graphics/filters/FilterEffect.cpp:
(WebCore):
(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.
(WebCore::FilterEffect::clearResultsRecursive):
(WebCore::FilterEffect::openCLImageToImageBuffer):
(WebCore::FilterEffect::createOpenCLImageResult):
(WebCore::FilterEffect::transformResultColorSpace):
* platform/graphics/filters/FilterEffect.h:
(FilterEffect):
(WebCore::FilterEffect::applyAll):
* platform/graphics/gpu/opencl/FilterContextOpenCL.cpp:
(WebCore::FilterContextOpenCL::isFailed):
(WebCore):
(WebCore::FilterContextOpenCL::freeResources):
(WebCore::FilterContextOpenCL::destroyContext):
(WebCore::FilterContextOpenCL::compileTransformColorSpaceProgram):
(WebCore::FilterContextOpenCL::openCLTransformColorSpace):
(WebCore::FilterContextOpenCL::compileProgram):
(WebCore::FilterContextOpenCL::freeResource):
* platform/graphics/gpu/opencl/FilterContextOpenCL.h:
(WebCore::FilterContextOpenCL::FilterContextOpenCL):
(WebCore::FilterContextOpenCL::setInError):
(WebCore::FilterContextOpenCL::inError):
(FilterContextOpenCL):
(WebCore::FilterContextOpenCL::RunKernel::RunKernel):
(WebCore::FilterContextOpenCL::RunKernel::addArgument):
(WebCore::FilterContextOpenCL::RunKernel::run):
(RunKernel):
* platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp:
(WebCore::FilterContextOpenCL::compileFEColorMatrix):
(WebCore::FEColorMatrix::platformApplyOpenCL):
* platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp:
(WebCore::FilterContextOpenCL::compileFETurbulence):
(WebCore::FETurbulence::platformApplyOpenCL):
* rendering/svg/RenderSVGResourceFilter.cpp:
(WebCore::RenderSVGResourceFilter::postApplyResource):
git-svn-id: https://svn.webkit.org/repository/webkit/trunk@142596
268f45cc-cd09-0410-ab3c-
d52691b4dbfc
+2013-02-12 Tamas Czene <tczene@inf.u-szeged.hu>
+
+ Add error checking into OpenCL version of SVG filters.
+ https://bugs.webkit.org/show_bug.cgi?id=107444
+
+ Reviewed by Zoltan Herczeg.
+
+ In case of an error the program runs through all the remaining filters by doing nothing.
+ After that deletes the results of every filter and starts software rendering.
+
+ * platform/graphics/filters/FilterEffect.cpp:
+ (WebCore):
+ (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.
+ (WebCore::FilterEffect::clearResultsRecursive):
+ (WebCore::FilterEffect::openCLImageToImageBuffer):
+ (WebCore::FilterEffect::createOpenCLImageResult):
+ (WebCore::FilterEffect::transformResultColorSpace):
+ * platform/graphics/filters/FilterEffect.h:
+ (FilterEffect):
+ (WebCore::FilterEffect::applyAll):
+ * platform/graphics/gpu/opencl/FilterContextOpenCL.cpp:
+ (WebCore::FilterContextOpenCL::isFailed):
+ (WebCore):
+ (WebCore::FilterContextOpenCL::freeResources):
+ (WebCore::FilterContextOpenCL::destroyContext):
+ (WebCore::FilterContextOpenCL::compileTransformColorSpaceProgram):
+ (WebCore::FilterContextOpenCL::openCLTransformColorSpace):
+ (WebCore::FilterContextOpenCL::compileProgram):
+ (WebCore::FilterContextOpenCL::freeResource):
+ * platform/graphics/gpu/opencl/FilterContextOpenCL.h:
+ (WebCore::FilterContextOpenCL::FilterContextOpenCL):
+ (WebCore::FilterContextOpenCL::setInError):
+ (WebCore::FilterContextOpenCL::inError):
+ (FilterContextOpenCL):
+ (WebCore::FilterContextOpenCL::RunKernel::RunKernel):
+ (WebCore::FilterContextOpenCL::RunKernel::addArgument):
+ (WebCore::FilterContextOpenCL::RunKernel::run):
+ (RunKernel):
+ * platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp:
+ (WebCore::FilterContextOpenCL::compileFEColorMatrix):
+ (WebCore::FEColorMatrix::platformApplyOpenCL):
+ * platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp:
+ (WebCore::FilterContextOpenCL::compileFETurbulence):
+ (WebCore::FETurbulence::platformApplyOpenCL):
+ * rendering/svg/RenderSVGResourceFilter.cpp:
+ (WebCore::RenderSVGResourceFilter::postApplyResource):
+
2013-02-12 Huang Dongsung <luxtella@company100.net>
[TexMap] Apply frames-per-second debug counter to WK1.
return m_inputEffects.at(number).get();
}
+#if ENABLE(OPENCL)
+void FilterEffect::applyAll()
+{
+ if (hasResult())
+ return;
+ FilterContextOpenCL* context = FilterContextOpenCL::context();
+ if (context) {
+ apply();
+ if (!context->inError())
+ return;
+ clearResultsRecursive();
+ context->destroyContext();
+ }
+ // Software code path.
+ apply();
+}
+#endif
+
void FilterEffect::apply()
{
if (hasResult())
#endif
}
+void FilterEffect::clearResultsRecursive()
+{
+ // Clear all results, regardless that the current effect has
+ // a result. Can be used if an effect is in an erroneous state.
+ if (hasResult())
+ clearResult();
+
+ unsigned size = m_inputEffects.size();
+ for (unsigned i = 0; i < size; ++i)
+ m_inputEffects.at(i).get()->clearResultsRecursive();
+}
+
ImageBuffer* FilterEffect::asImageBuffer()
{
if (!hasResult())
FilterContextOpenCL* context = FilterContextOpenCL::context();
ASSERT(context);
+ if (context->inError())
+ return 0;
+
size_t origin[3] = { 0, 0, 0 };
size_t region[3] = { m_absolutePaintRect.width(), m_absolutePaintRect.height(), 1 };
RefPtr<Uint8ClampedArray> destinationPixelArray = Uint8ClampedArray::create(m_absolutePaintRect.width() * m_absolutePaintRect.height() * 4);
- clFinish(context->commandQueue());
- clEnqueueReadImage(context->commandQueue(), m_openCLImageResult, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0);
+ if (context->isFailed(clFinish(context->commandQueue())))
+ return 0;
+
+ if (context->isFailed(clEnqueueReadImage(context->commandQueue(), m_openCLImageResult, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0)))
+ return 0;
m_imageBufferResult = ImageBuffer::create(m_absolutePaintRect.size());
IntRect destinationRect(IntPoint(), m_absolutePaintRect.size());
#if ENABLE(OPENCL)
OpenCLHandle FilterEffect::createOpenCLImageResult(uint8_t* source)
{
+ FilterContextOpenCL* context = FilterContextOpenCL::context();
+ ASSERT(context);
+
+ if (context->inError())
+ return 0;
+
ASSERT(!hasResult());
cl_image_format clImageFormat;
clImageFormat.image_channel_order = CL_RGBA;
clImageFormat.image_channel_data_type = CL_UNORM_INT8;
- FilterContextOpenCL* context = FilterContextOpenCL::context();
- ASSERT(context);
+ int errorCode = 0;
m_openCLImageResult = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE | (source ? CL_MEM_COPY_HOST_PTR : 0),
- &clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, 0);
+ &clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, &errorCode);
+ if (context->isFailed(errorCode))
+ return 0;
+
return m_openCLImageResult;
}
#endif
#if ENABLE(OPENCL)
if (openCLImage()) {
+ if (m_imageBufferResult)
+ m_imageBufferResult.clear();
FilterContextOpenCL* context = FilterContextOpenCL::context();
ASSERT(context);
context->openCLTransformColorSpace(m_openCLImageResult, absolutePaintRect(), m_resultColorSpace, dstColorSpace);
- if (m_imageBufferResult)
- m_imageBufferResult.clear();
} else {
#endif
- // FIXME: We can avoid this potentially unnecessary ImageBuffer conversion by adding
- // color space transform support for the {pre,un}multiplied arrays.
- asImageBuffer()->transformColorSpace(m_resultColorSpace, dstColorSpace);
+ // FIXME: We can avoid this potentially unnecessary ImageBuffer conversion by adding
+ // color space transform support for the {pre,un}multiplied arrays.
+ asImageBuffer()->transformColorSpace(m_resultColorSpace, dstColorSpace);
#if ENABLE(OPENCL)
}
virtual ~FilterEffect();
void clearResult();
+ void clearResultsRecursive();
+
ImageBuffer* asImageBuffer();
PassRefPtr<Uint8ClampedArray> asUnmultipliedImage(const IntRect&);
PassRefPtr<Uint8ClampedArray> asPremultipliedImage(const IntRect&);
void setMaxEffectRect(const FloatRect& maxEffectRect) { m_maxEffectRect = maxEffectRect; }
void apply();
-
+#if ENABLE(OPENCL)
+ void applyAll();
+#else
+ inline void applyAll() { apply(); }
+#endif
+
// Correct any invalid pixels, if necessary, in the result of a filter operation.
// This method is used to ensure valid pixel values on filter inputs and the final result.
// Only the arithmetic composite filter ever needs to perform correction.
return m_context;
}
+void FilterContextOpenCL::freeResources()
+{
+ clFinish(m_commandQueue);
+
+ if (m_colorMatrixWasCompiled) {
+ freeResource(m_matrixOperation);
+ freeResource(m_saturateAndHueRotateOperation);
+ freeResource(m_luminanceOperation);
+ freeResource(m_colorMatrixProgram);
+ }
+ m_colorMatrixWasCompiled = false;
+
+ if (m_turbulenceWasCompiled) {
+ freeResource(m_turbulenceOperation);
+ freeResource(m_turbulenceProgram);
+ }
+ m_turbulenceWasCompiled = false;
+
+ if (m_transformColorSpaceWasCompiled) {
+ freeResource(m_transformColorSpaceKernel);
+ freeResource(m_transformColorSpaceProgram);
+ }
+ m_transformColorSpaceWasCompiled = false;
+}
+
+void FilterContextOpenCL::destroyContext()
+{
+ freeResources();
+
+ if (m_commandQueue)
+ clReleaseCommandQueue(m_commandQueue);
+ m_commandQueue = 0;
+
+ if (m_deviceContext)
+ clReleaseContext(m_deviceContext);
+ m_deviceContext = 0;
+
+ m_context = 0;
+}
+
OpenCLHandle FilterContextOpenCL::createOpenCLImage(IntSize paintSize)
{
FilterContextOpenCL* context = FilterContextOpenCL::context();
}
); // End of OpenCL kernels
+inline bool FilterContextOpenCL::compileTransformColorSpaceProgram()
+{
+ if (m_transformColorSpaceWasCompiled || inError())
+ return !inError();
+
+ m_transformColorSpaceWasCompiled = true;
+
+ if (isResourceAllocationFailed((m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram))))
+ return false;
+ if (isResourceAllocationFailed((m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace"))))
+ return false;
+ return true;
+}
+
void FilterContextOpenCL::openCLTransformColorSpace(OpenCLHandle& source, IntRect sourceSize, ColorSpace srcColorSpace, ColorSpace dstColorSpace)
{
DEFINE_STATIC_LOCAL(OpenCLHandle, deviceRgbLUT, ());
DEFINE_STATIC_LOCAL(OpenCLHandle, linearRgbLUT, ());
- if (srcColorSpace == dstColorSpace)
+ if (srcColorSpace == dstColorSpace || inError())
return;
if ((srcColorSpace != ColorSpaceLinearRGB && srcColorSpace != ColorSpaceDeviceRGB)
|| (dstColorSpace != ColorSpaceLinearRGB && dstColorSpace != ColorSpaceDeviceRGB))
return;
- FilterContextOpenCL* context = FilterContextOpenCL::context();
- ASSERT(context);
+ if (!compileTransformColorSpaceProgram())
+ return;
- OpenCLHandle destination = context->createOpenCLImage(sourceSize.size());
+ OpenCLHandle destination = createOpenCLImage(sourceSize.size());
- if (!m_transformColorSpaceProgram) {
- m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram);
- ASSERT(m_transformColorSpaceProgram);
- m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace");
- ASSERT(m_transformColorSpaceKernel);
- }
-
- RunKernel kernel(context, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height());
+ RunKernel kernel(this, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height());
kernel.addArgument(source);
kernel.addArgument(destination);
cl_program FilterContextOpenCL::compileProgram(const char* source)
{
cl_program program;
- cl_int errorNumber;
+ cl_int errorNumber = 0;
- FilterContextOpenCL* context = FilterContextOpenCL::context();
- ASSERT(context);
+ program = clCreateProgramWithSource(m_deviceContext, 1, (const char**) &source, 0, &errorNumber);
+ if (isFailed(errorNumber))
+ return 0;
- program = clCreateProgramWithSource(context->m_deviceContext, 1, (const char**) &source, 0, 0);
- errorNumber = clBuildProgram(program, 0, 0, 0, 0, 0);
- if (errorNumber)
+ if (isFailed(clBuildProgram(program, 0, 0, 0, 0, 0)))
return 0;
return program;
}
+
+void FilterContextOpenCL::freeResource(cl_kernel& handle)
+{
+ if (handle) {
+ clReleaseKernel(handle);
+ handle = 0;
+ }
+}
+
+void FilterContextOpenCL::freeResource(cl_program& handle)
+{
+ if (handle) {
+ clReleaseProgram(handle);
+ handle = 0;
+ }
+}
} // namespace WebCore
#endif
class FilterContextOpenCL {
public:
FilterContextOpenCL()
- : m_deviceId(0)
+ : m_inError(false)
+ , m_deviceId(0)
, m_deviceContext(0)
, m_commandQueue(0)
+ , m_transformColorSpaceWasCompiled(false)
, m_transformColorSpaceProgram(0)
, m_transformColorSpaceKernel(0)
- , m_colorMatrixCompileStatus(openclNotCompiledYet)
+ , m_colorMatrixWasCompiled(false)
, m_colorMatrixProgram(0)
, m_matrixOperation(0)
, m_saturateAndHueRotateOperation(0)
, m_luminanceOperation(0)
- , m_turbulenceCompileStatus(openclNotCompiledYet)
+ , m_turbulenceWasCompiled(false)
, m_turbulenceProgram(0)
, m_turbulenceOperation(0)
{
cl_context deviceContext() { return m_deviceContext; }
cl_command_queue commandQueue() { return m_commandQueue; }
+ inline void setInError(bool errorCode = true) { m_inError = errorCode; }
+ inline bool inError() { return m_inError; }
+ inline bool isFailed(bool);
+ inline bool isResourceAllocationFailed(bool);
+
+ void freeResources();
+ void destroyContext();
+
OpenCLHandle createOpenCLImage(IntSize);
+
+ inline bool compileTransformColorSpaceProgram();
void openCLTransformColorSpace(OpenCLHandle&, IntRect, ColorSpace, ColorSpace);
inline bool compileFEColorMatrix();
RunKernel(FilterContextOpenCL* context, cl_kernel kernel, size_t width, size_t height)
: m_context(context)
, m_kernel(kernel)
- , index(0)
+ , m_index(0)
+ , m_error(context->inError())
{
m_globalSize[0] = width;
m_globalSize[1] = height;
void addArgument(OpenCLHandle handle)
{
- clSetKernelArg(m_kernel, index++, sizeof(OpenCLHandle), handle.handleAddress());
+ if (!m_error)
+ m_error = clSetKernelArg(m_kernel, m_index++, sizeof(OpenCLHandle), handle.handleAddress());
}
void addArgument(cl_int value)
{
- clSetKernelArg(m_kernel, index++, sizeof(cl_int), reinterpret_cast<void*>(&value));
+ if (!m_error)
+ m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_int), reinterpret_cast<void*>(&value));
}
void addArgument(cl_float value)
{
- clSetKernelArg(m_kernel, index++, sizeof(cl_float), reinterpret_cast<void*>(&value));
+ if (!m_error)
+ m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_float), reinterpret_cast<void*>(&value));
}
void addArgument(cl_sampler handle)
{
- clSetKernelArg(m_kernel, index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle));
+ if (!m_error)
+ m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle));
}
OpenCLHandle addArgument(void* buffer, int size)
{
- OpenCLHandle handle(clCreateBuffer(m_context->deviceContext(), CL_MEM_READ_ONLY, size, 0, 0));
- clEnqueueWriteBuffer(m_context->commandQueue(), handle, CL_TRUE, 0, size, buffer, 0, 0, 0);
- clSetKernelArg(m_kernel, index++, sizeof(OpenCLHandle), handle.handleAddress());
- return handle;
+ if (m_error)
+ return 0;
+ OpenCLHandle handle(clCreateBuffer(m_context->deviceContext(), CL_MEM_READ_ONLY, size, 0, &m_error));
+ if (m_error)
+ return 0;
+ m_error = clEnqueueWriteBuffer(m_context->commandQueue(), handle, CL_TRUE, 0, size, buffer, 0, 0, 0);
+ if (m_error)
+ return 0;
+ m_error = clSetKernelArg(m_kernel, m_index++, sizeof(OpenCLHandle), handle.handleAddress());
+ return !m_error ? handle : 0;
}
void run()
{
- clFinish(m_context->m_commandQueue);
- clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0);
+ if (m_context->isFailed(m_error))
+ return;
+
+ m_error = clFinish(m_context->m_commandQueue);
+ if (!m_error)
+ m_error = clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0);
+ m_context->isFailed(m_error);
}
FilterContextOpenCL* m_context;
cl_kernel m_kernel;
size_t m_globalSize[2];
- int index;
+ int m_index;
+ int m_error;
};
- enum OpenCLCompileStatus {
- openclNotCompiledYet,
- openclCompileFailed,
- openclCompileSuccessful
- };
-
- static cl_program compileProgram(const char*);
+ cl_program compileProgram(const char*);
static inline cl_kernel kernelByName(cl_program program, const char* name) { return clCreateKernel(program, name, 0); }
+ static inline void freeResource(cl_kernel&);
+ static inline void freeResource(cl_program&);
+
static FilterContextOpenCL* m_context;
static int m_alreadyInitialized;
+ bool m_inError;
cl_device_id m_deviceId;
cl_context m_deviceContext;
cl_command_queue m_commandQueue;
+ bool m_transformColorSpaceWasCompiled;
cl_program m_transformColorSpaceProgram;
cl_kernel m_transformColorSpaceKernel;
- OpenCLCompileStatus m_colorMatrixCompileStatus;
+ bool m_colorMatrixWasCompiled;
cl_program m_colorMatrixProgram;
cl_kernel m_matrixOperation;
cl_kernel m_saturateAndHueRotateOperation;
cl_kernel m_luminanceOperation;
- OpenCLCompileStatus m_turbulenceCompileStatus;
+ bool m_turbulenceWasCompiled;
cl_program m_turbulenceProgram;
cl_kernel m_turbulenceOperation;
};
+inline bool FilterContextOpenCL::isFailed(bool value)
+{
+ if (value)
+ setInError();
+ return value;
+}
+
+inline bool FilterContextOpenCL::isResourceAllocationFailed(bool value)
+{
+ if (!value)
+ setInError();
+ return !value;
+}
+
} // namespace WebCore
#endif // ENABLE(OPENCL)
inline bool FilterContextOpenCL::compileFEColorMatrix()
{
- if (m_colorMatrixCompileStatus != openclNotCompiledYet)
- return m_colorMatrixCompileStatus == openclCompileSuccessful;
+ if (m_colorMatrixWasCompiled || inError())
+ return !inError();
- m_colorMatrixCompileStatus = openclCompileFailed;
- m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram);
- if (!m_colorMatrixProgram)
- return false;
+ m_colorMatrixWasCompiled = true;
- m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix");
- if (!m_matrixOperation)
+ if (isResourceAllocationFailed((m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram))))
return false;
- m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate");
- if (!m_saturateAndHueRotateOperation)
+ if (isResourceAllocationFailed((m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix"))))
return false;
- m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance");
- if (!m_luminanceOperation)
+ if (isResourceAllocationFailed((m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate"))))
return false;
-
- m_colorMatrixCompileStatus = openclCompileSuccessful;
- return openclCompileSuccessful;
+ if (isResourceAllocationFailed((m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate"))))
+ return false;
+ if (isResourceAllocationFailed((m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance"))))
+ return false;
+ return true;
}
inline void FilterContextOpenCL::applyFEColorMatrix(OpenCLHandle destination, IntSize destinationSize, OpenCLHandle source, IntPoint relativeSourceLocation, float* values, int type)
bool FEColorMatrix::platformApplyOpenCL()
{
FilterContextOpenCL* context = FilterContextOpenCL::context();
- if (!context || !context->compileFEColorMatrix())
+ if (!context)
return false;
+ if (!context->compileFEColorMatrix())
+ return true;
+
FilterEffect* in = inputEffect(0);
OpenCLHandle source = in->openCLImage();
OpenCLHandle destination = createOpenCLImageResult();
inline bool FilterContextOpenCL::compileFETurbulence()
{
- if (m_turbulenceCompileStatus != openclNotCompiledYet)
- return m_turbulenceCompileStatus == openclCompileSuccessful;
+ if (m_turbulenceWasCompiled || inError())
+ return !inError();
- m_turbulenceCompileStatus = openclCompileFailed;
- m_turbulenceProgram = compileProgram(turbulenceKernelProgram);
- if (!m_turbulenceProgram)
+ m_turbulenceWasCompiled = true;
+
+ if (isResourceAllocationFailed((m_turbulenceProgram = compileProgram(turbulenceKernelProgram))))
return false;
- m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence");
- if (!m_turbulenceOperation)
+ if (isResourceAllocationFailed((m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence"))))
return false;
-
- m_turbulenceCompileStatus = openclCompileSuccessful;
- return openclCompileSuccessful;
+ return true;
}
inline void FilterContextOpenCL::applyFETurbulence(OpenCLHandle destination,
bool FETurbulence::platformApplyOpenCL()
{
FilterContextOpenCL* context = FilterContextOpenCL::context();
- if (!context || !context->compileFETurbulence())
+ if (!context)
return false;
+ if (!context->compileFETurbulence())
+ return true;
+
OpenCLHandle destination = createOpenCLImageResult();
PaintingData paintingData(m_seed, roundedIntSize(filterPrimitiveSubregion().size()));
// Always true if filterData is just built (filterData->state == FilterData::Built).
if (!lastEffect->hasResult()) {
filterData->state = FilterData::Applying;
- lastEffect->apply();
+ lastEffect->applyAll();
lastEffect->correctFilterResultIfNeeded();
lastEffect->transformResultColorSpace(ColorSpaceDeviceRGB);
}
git://git.webkit.org / WebKit-https.git / commitdiff