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; |
-} |