| Index: experimental/skpdiff/SkDifferentPixelsMetric_opencl.cpp
|
| diff --git a/experimental/skpdiff/SkDifferentPixelsMetric_opencl.cpp b/experimental/skpdiff/SkDifferentPixelsMetric_opencl.cpp
|
| deleted file mode 100644
|
| index 61fb63813bcabbbcedbe5575796c3876f24573e1..0000000000000000000000000000000000000000
|
| --- a/experimental/skpdiff/SkDifferentPixelsMetric_opencl.cpp
|
| +++ /dev/null
|
| @@ -1,168 +0,0 @@
|
| -/*
|
| - * Copyright 2013 Google Inc.
|
| - *
|
| - * Use of this source code is governed by a BSD-style license that can be
|
| - * found in the LICENSE file.
|
| - */
|
| -
|
| -#include "SkBitmap.h"
|
| -
|
| -#include "SkDifferentPixelsMetric.h"
|
| -#include "skpdiff_util.h"
|
| -
|
| -static const char kDifferentPixelsKernelSource[] =
|
| - "#pragma OPENCL_EXTENSION cl_khr_global_int32_base_atomics \n"
|
| - " \n"
|
| - "const sampler_t gInSampler = CLK_NORMALIZED_COORDS_FALSE | \n"
|
| - " CLK_ADDRESS_CLAMP_TO_EDGE | \n"
|
| - " CLK_FILTER_NEAREST; \n"
|
| - " \n"
|
| - "__kernel void diff(read_only image2d_t baseline, read_only image2d_t test, \n"
|
| - " __global int* result, __global int2* poi) { \n"
|
| - " int2 coord = (int2)(get_global_id(0), get_global_id(1)); \n"
|
| - " uint4 baselinePixel = read_imageui(baseline, gInSampler, coord); \n"
|
| - " uint4 testPixel = read_imageui(test, gInSampler, coord); \n"
|
| - " int4 pixelCompare = baselinePixel == testPixel; \n"
|
| - " if (baselinePixel.x != testPixel.x || \n"
|
| - " baselinePixel.y != testPixel.y || \n"
|
| - " baselinePixel.z != testPixel.z || \n"
|
| - " baselinePixel.w != testPixel.w) { \n"
|
| - " \n"
|
| - " int poiIndex = atomic_inc(result); \n"
|
| - " poi[poiIndex] = coord; \n"
|
| - " } \n"
|
| - "} \n";
|
| -
|
| -struct SkDifferentPixelsMetric::QueuedDiff {
|
| - bool finished;
|
| - double result;
|
| - int numDiffPixels;
|
| - SkIPoint* poi;
|
| - cl_mem baseline;
|
| - cl_mem test;
|
| - cl_mem resultsBuffer;
|
| - cl_mem poiBuffer;
|
| -};
|
| -
|
| -const char* SkDifferentPixelsMetric::getName() {
|
| - return "different_pixels";
|
| -}
|
| -
|
| -int SkDifferentPixelsMetric::queueDiff(SkBitmap* baseline, SkBitmap* test) {
|
| - int diffID = fQueuedDiffs.count();
|
| - double startTime = get_seconds();
|
| - QueuedDiff* diff = fQueuedDiffs.push();
|
| -
|
| - // If we never end up running the kernel, include some safe defaults in the result.
|
| - diff->finished = false;
|
| - diff->result = -1.0;
|
| - diff->numDiffPixels = 0;
|
| - diff->poi = NULL;
|
| -
|
| - // Ensure the images are comparable
|
| - if (baseline->width() != test->width() || baseline->height() != test->height() ||
|
| - baseline->width() <= 0 || baseline->height() <= 0 ||
|
| - baseline->config() != test->config()) {
|
| - diff->finished = true;
|
| - return diffID;
|
| - }
|
| -
|
| - // Upload images to the CL device
|
| - if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test, &diff->test)) {
|
| - diff->finished = true;
|
| - fIsGood = false;
|
| - return -1;
|
| - }
|
| -
|
| - // A small hack that makes calculating percentage difference easier later on.
|
| - diff->result = 1.0 / ((double)baseline->width() * baseline->height());
|
| -
|
| - // Make a buffer to store results into. It must be initialized with pointers to memory.
|
| - static const int kZero = 0;
|
| - // We know OpenCL won't write to it because we use CL_MEM_COPY_HOST_PTR
|
| - diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
|
| - sizeof(int), (int*)&kZero, NULL);
|
| -
|
| - diff->poiBuffer = clCreateBuffer(fContext, CL_MEM_WRITE_ONLY,
|
| - sizeof(int) * 2 * baseline->width() * baseline->height(),
|
| - NULL, NULL);
|
| -
|
| - // Set all kernel arguments
|
| - cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baseline);
|
| - setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test);
|
| - setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->resultsBuffer);
|
| - setArgErr |= clSetKernelArg(fKernel, 3, sizeof(cl_mem), &diff->poiBuffer);
|
| - if (CL_SUCCESS != setArgErr) {
|
| - SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr));
|
| - fIsGood = false;
|
| - return -1;
|
| - }
|
| -
|
| - // Queue this diff on the CL device
|
| - cl_event event;
|
| - const size_t workSize[] = { baseline->width(), baseline->height() };
|
| - cl_int enqueueErr;
|
| - enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSize,
|
| - NULL, 0, NULL, &event);
|
| - if (CL_SUCCESS != enqueueErr) {
|
| - SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr));
|
| - fIsGood = false;
|
| - return -1;
|
| - }
|
| -
|
| - // This makes things totally synchronous. Actual queue is not ready yet
|
| - clWaitForEvents(1, &event);
|
| - diff->finished = true;
|
| -
|
| - // Immediate read back the results
|
| - clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0,
|
| - sizeof(int), &diff->numDiffPixels, 0, NULL, NULL);
|
| - diff->result *= (double)diff->numDiffPixels;
|
| - diff->result = (1.0 - diff->result);
|
| -
|
| - diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels);
|
| - clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0,
|
| - sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NULL, NULL);
|
| -
|
| - // Release all the buffers created
|
| - clReleaseMemObject(diff->poiBuffer);
|
| - clReleaseMemObject(diff->resultsBuffer);
|
| - clReleaseMemObject(diff->baseline);
|
| - clReleaseMemObject(diff->test);
|
| -
|
| - SkDebugf("Time: %f\n", (get_seconds() - startTime));
|
| -
|
| - return diffID;
|
| -}
|
| -
|
| -void SkDifferentPixelsMetric::deleteDiff(int id) {
|
| - QueuedDiff* diff = &fQueuedDiffs[id];
|
| - if (NULL != diff->poi) {
|
| - SkDELETE_ARRAY(diff->poi);
|
| - diff->poi = NULL;
|
| - }
|
| -}
|
| -
|
| -bool SkDifferentPixelsMetric::isFinished(int id) {
|
| - return fQueuedDiffs[id].finished;
|
| -}
|
| -
|
| -double SkDifferentPixelsMetric::getResult(int id) {
|
| - return fQueuedDiffs[id].result;
|
| -}
|
| -
|
| -int SkDifferentPixelsMetric::getPointsOfInterestCount(int id) {
|
| - return fQueuedDiffs[id].numDiffPixels;
|
| -}
|
| -
|
| -SkIPoint* SkDifferentPixelsMetric::getPointsOfInterest(int id) {
|
| - return fQueuedDiffs[id].poi;
|
| -}
|
| -
|
| -bool SkDifferentPixelsMetric::onInit() {
|
| - if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel)) {
|
| - return false;
|
| - }
|
| -
|
| - return true;
|
| -}
|
|
|