add skpdiff tool to compare bitmaps

- start framework for pluggable algorithms
- implement simple number of pixels different OpenCL algo

R=djsollen@google.com, bsalomon@google.com, jvanverth@google.com

Author: zachr@google.com

Review URL: https://chromiumcodereview.appspot.com/16284007

git-svn-id: http://skia.googlecode.com/svn/trunk@9616 2bbb7eff-a529-9590-31e7-b0007b416f81

Makefile

@@ -48,7 +48,8 @@ VALID_TARGETS := \
                  SkiaAndroidApp \
                  skia_lib \
                  tests \
-                 tools
+                 tools \
+                 skpdiff
 
 # Default target.  This must be listed before all other targets.
 .PHONY: default

experimental/skpdiff/SkCLImageDiffer.cpp

@@ -0,0 +1,210 @@
+
+/*
+ * 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 <cstring>
+
+#include "SkBitmap.h"
+#include "SkStream.h"
+
+#include "SkCLImageDiffer.h"
+#include "skpdiff_util.h"
+
+SkCLImageDiffer::SkCLImageDiffer() {
+    fIsGood = false;
+}
+
+
+bool SkCLImageDiffer::init(cl_device_id device, cl_context context) {
+    fContext = context;
+    fDevice = device;
+
+    cl_int queueErr;
+    fCommandQueue = clCreateCommandQueue(fContext, fDevice, 0, &queueErr);
+    if (CL_SUCCESS != queueErr) {
+        SkDebugf("Command queue creation failed: %s\n", cl_error_to_string(queueErr));
+        fIsGood = false;
+        return false;
+    }
+
+    fIsGood = this->onInit();
+    return fIsGood;
+}
+
+bool SkCLImageDiffer::loadKernelFile(const char file[], const char name[], cl_kernel* kernel) {
+    // Open the kernel source file
+    SkFILEStream sourceStream(file);
+    if (!sourceStream.isValid()) {
+        SkDebugf("Failed to open kernel source file");
+        return false;
+    }
+
+    return loadKernelStream(&sourceStream, name, kernel);
+}
+
+bool SkCLImageDiffer::loadKernelStream(SkStream* stream, const char name[], cl_kernel* kernel) {
+    // Read the kernel source into memory
+    SkString sourceString;
+    sourceString.resize(stream->getLength());
+    size_t bytesRead = stream->read(sourceString.writable_str(), sourceString.size());
+    if (bytesRead != sourceString.size()) {
+        SkDebugf("Failed to read kernel source file");
+        return false;
+    }
+
+    return loadKernelSource(sourceString.c_str(), name, kernel);
+}
+
+bool SkCLImageDiffer::loadKernelSource(const char source[], const char name[], cl_kernel* kernel) {
+    // Build the kernel source
+    size_t sourceLen = strlen(source);
+    cl_program program = clCreateProgramWithSource(fContext, 1, &source, &sourceLen, NULL);
+    cl_int programErr = clBuildProgram(program, 1, &fDevice, "", NULL, NULL);
+    if (CL_SUCCESS != programErr) {
+        SkDebugf("Program creation failed: %s\n", cl_error_to_string(programErr));
+
+        // Attempt to get information about why the build failed
+        char buildLog[4096];
+        clGetProgramBuildInfo(program, fDevice, CL_PROGRAM_BUILD_LOG, sizeof(buildLog), buildLog, NULL);
+        SkDebugf("Build log: %s\n", buildLog);
+
+        return false;
+    }
+
+    cl_int kernelErr;
+    *kernel = clCreateKernel(program, name, &kernelErr);
+    if (CL_SUCCESS != kernelErr) {
+        SkDebugf("Kernel creation failed: %s\n", cl_error_to_string(kernelErr));
+        return false;
+    }
+
+    return true;
+}
+
+bool SkCLImageDiffer::makeImage2D(SkBitmap* bitmap, cl_mem* image) {
+    cl_int imageErr;
+    cl_image_format bitmapFormat;
+    switch (bitmap->config()) {
+        case SkBitmap::kA8_Config:
+            bitmapFormat.image_channel_order = CL_A;
+            bitmapFormat.image_channel_data_type = CL_UNSIGNED_INT8;
+            break;
+        case SkBitmap::kRGB_565_Config:
+            bitmapFormat.image_channel_order = CL_RGB;
+            bitmapFormat.image_channel_data_type = CL_UNORM_SHORT_565;
+            break;
+        case SkBitmap::kARGB_8888_Config:
+            bitmapFormat.image_channel_order = CL_RGBA;
+            bitmapFormat.image_channel_data_type = CL_UNSIGNED_INT8;
+            break;
+        default:
+            SkDebugf("Image format is unsupported\n");
+            return false;
+    }
+
+    // Upload the bitmap data to OpenCL
+    bitmap->lockPixels();
+    *image = clCreateImage2D(fContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
+                             &bitmapFormat, bitmap->width(), bitmap->height(),
+                             bitmap->rowBytes(), bitmap->getPixels(),
+                             &imageErr);
+    bitmap->unlockPixels();
+
+    if (CL_SUCCESS != imageErr) {
+        SkDebugf("Input image creation failed: %s\n", cl_error_to_string(imageErr));
+        return false;
+    }
+
+    return true;
+}
+
+
+////////////////////////////////////////////////////////////////
+
+const char* SkDifferentPixelsImageDiffer::getName() {
+    return "Find Different Pixels";
+}
+
+int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test) {
+    int diffID = fQueuedDiffs.count();
+    double startTime = get_seconds();
+    QueuedDiff* diff = fQueuedDiffs.push();
+
+    // Ensure the images are comparable
+    if (baseline->width() != test->width() || baseline->height() != test->height() ||
+                    baseline->width() <= 0 || baseline->height() <= 0) {
+        diff->finished = true;
+        diff->result = 0.0;
+        return diffID;
+    }
+
+    // Upload images to the CL device
+    if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test, &diff->test)) {
+        diff->finished = true;
+        diff->result = 0.0;
+        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
+    int numDiffPixels = 0;
+    diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
+                                         sizeof(int), &numDiffPixels, 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);
+    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), &numDiffPixels, 0, NULL, NULL);
+    diff->result *= (double)numDiffPixels;
+    diff->result = (1.0 - diff->result);
+    SkDebugf("Time: %f\n", (get_seconds() - startTime));
+
+    return diffID;
+}
+
+bool SkDifferentPixelsImageDiffer::isFinished(int id) {
+    return fQueuedDiffs[id].finished;
+}
+
+double SkDifferentPixelsImageDiffer::getResult(int id) {
+    return fQueuedDiffs[id].result;
+}
+
+
+bool SkDifferentPixelsImageDiffer::onInit() {
+    if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel)) {
+        return false;
+    }
+
+    return true;
+}

experimental/skpdiff/SkCLImageDiffer.h

@@ -0,0 +1,117 @@
+/*
+ * Copyright 2013 Google Inc.
+ *
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#ifndef SkCLImageDiffer_DEFINED
+#define SkCLImageDiffer_DEFINED
+
+#include <CL/cl.h>
+#include "SkTDArray.h"
+
+#include "SkImageDiffer.h"
+
+class SkStream;
+
+/**
+ * An SkImageDiffer that requires initialization with an OpenCL device and context.
+ */
+class SkCLImageDiffer : public SkImageDiffer {
+public:
+    SkCLImageDiffer();
+
+    /**
+     * Initializes the OpenCL resources this differ needs to work
+     * @param  device  An OpenCL device
+     * @param  context An OpenCL context of the given device
+     * @return         True on success, false otherwise
+     */
+    virtual bool init(cl_device_id device, cl_context context);
+
+protected:
+    /**
+     * Called by init after fDevice, fContext, and fCommandQueue are successfully initialized
+     * @return True on success, false otherwise
+     */
+    virtual bool onInit() = 0;
+
+    /**
+     * Loads an OpenCL kernel from the file with the given named entry point. This only works after
+     * init is called.
+     * @param  file   The file path of the kernel
+     * @param  name   The name of the entry point of the desired kernel in the file
+     * @param  kernel A pointer to return the loaded kernel into
+     * @return        True on success, false otherwise
+     */
+    bool loadKernelFile(const char file[], const char name[], cl_kernel* kernel);
+
+    /**
+     * Loads an OpenCL kernel from the stream with the given named entry point. This only works
+     * after init is called.
+     * @param  stream  The stream that contains the kernel
+     * @param  name    The name of the entry point of the desired kernel in the stream
+     * @param  kernel  A pointer to return the loaded kernel into
+     * @return         True on success, false otherwise
+     */
+    bool loadKernelStream(SkStream* stream, const char name[], cl_kernel* kernel);
+
+    /**
+     * Loads an OpenCL kernel from the source string with the given named entry point. This only
+     * works after init is called.
+     * @param  source  The string that contains the kernel
+     * @param  name    The name of the entry point of the desired kernel in the source string
+     * @param  kernel  A pointer to return the loaded kernel into
+     * @return         True on success, false otherwise
+     */
+    bool loadKernelSource(const char source[], const char name[], cl_kernel* kernel);
+
+    /**
+     * Loads a read only copy of the given bitmap into device memory and returns the block of
+     * memory. This only works after init is called.
+     * @param  bitmap The bitmap to load into memory
+     * @param  image  A pointer to return the allocated image to
+     * @return        True on success, false otherwise
+     */
+    bool makeImage2D(SkBitmap* bitmap, cl_mem* image);
+
+    cl_device_id     fDevice;
+    cl_context       fContext;
+    cl_command_queue fCommandQueue;
+
+private:
+
+    typedef SkImageDiffer INHERITED;
+};
+
+/**
+ * A OpenCL differ that measures the percentage of different corresponding pixels. If the two images
+ * are not the same size or have no pixels, the result will always be zero.
+ */
+class SkDifferentPixelsImageDiffer : public SkCLImageDiffer {
+public:
+    virtual const char* getName() SK_OVERRIDE;
+    virtual int queueDiff(SkBitmap* baseline, SkBitmap* test) SK_OVERRIDE;
+    virtual bool isFinished(int id) SK_OVERRIDE;
+    virtual double getResult(int id) SK_OVERRIDE;
+
+protected:
+    virtual bool onInit() SK_OVERRIDE;
+
+private:
+    struct QueuedDiff {
+        bool finished;
+        double result;
+        cl_mem baseline;
+        cl_mem test;
+        cl_mem resultsBuffer;
+    };
+
+    SkTDArray<QueuedDiff> fQueuedDiffs;
+    cl_kernel fKernel;
+
+    typedef SkCLImageDiffer INHERITED;
+};
+
+#endif

experimental/skpdiff/SkImageDiffer.cpp

@@ -0,0 +1,39 @@
+/*
+ * 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 <cstring>
+
+#include "SkBitmap.h"
+#include "SkImageDecoder.h"
+
+#include "SkImageDiffer.h"
+
+#include "skpdiff_util.h"
+
+
+SkImageDiffer::SkImageDiffer()
+    : fIsGood(true) {
+
+}
+
+SkImageDiffer::~SkImageDiffer() {
+
+}
+
+int SkImageDiffer::queueDiffOfFile(const char baseline[], const char test[]) {
+    SkBitmap baselineBitmap;
+    SkBitmap testBitmap;
+    if (!SkImageDecoder::DecodeFile(baseline, &baselineBitmap)) {
+        SkDebugf("Failed to load bitmap \"%s\"\n", baseline);
+        return -1;
+    }
+    if (!SkImageDecoder::DecodeFile(test, &testBitmap)) {
+        SkDebugf("Failed to load bitmap \"%s\"\n", test);
+        return -1;
+    }
+    return this->queueDiff(&baselineBitmap, &testBitmap);
+}

experimental/skpdiff/SkImageDiffer.h

@@ -0,0 +1,70 @@
+/*
+ * Copyright 2013 Google Inc.
+ *
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#ifndef SkImageDiffer_DEFINED
+#define SkImageDiffer_DEFINED
+
+class SkBitmap;
+
+/**
+ * Encapsulates an image difference metric algorithm that can be potentially run asynchronously.
+ */
+class SkImageDiffer {
+public:
+    SkImageDiffer();
+    virtual ~SkImageDiffer();
+
+    /**
+     * Gets a unique and descriptive name of this differ
+     * @return A statically allocated null terminated string that is the name of this differ
+     */
+    virtual const char* getName() = 0;
+
+    /**
+     * Gets if this differ is in a usable state
+     * @return True if this differ can be used, false otherwise
+     */
+    bool isGood() { return fIsGood; }
+
+    /**
+     * Wraps a call to queueDiff by loading the given filenames into SkBitmaps
+     * @param  baseline The file path of the baseline image
+     * @param  test     The file path of the test image
+     * @return          The results of queueDiff with the loaded bitmaps
+     */
+    int queueDiffOfFile(const char baseline[], const char test[]);
+
+    /**
+     * Queues a diff on a pair of bitmaps to be done at some future time.
+     * @param  baseline The correct bitmap
+     * @param  test     The bitmap whose difference is being tested
+     * @return          An non-negative diff ID on success, a negative integer on failure.
+     */
+    virtual int queueDiff(SkBitmap* baseline, SkBitmap* test) = 0;
+
+    /**
+     * Gets whether a queued diff of the given id has finished
+     * @param  id The id of the queued diff to query
+     * @return    True if the queued diff is finished and has results, false otherwise
+     */
+    virtual bool isFinished(int id) = 0;
+
+    /**
+     * Gets the results of the queued diff of the given id. The results are only meaningful after
+     * the queued diff has finished.
+     * @param  id The id of the queued diff to query
+     * @return    A score between of how different the compared images are, with lower numbers being
+     *            more different.
+     */
+    virtual double getResult(int id) = 0;
+
+protected:
+    bool fIsGood;
+};
+
+
+#endif

experimental/skpdiff/diff_pixels.cl

@@ -0,0 +1,26 @@
+/*
+ * Copyright 2013 Google Inc.
+ *
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#pragma OPENCL_EXTENSION cl_khr_global_int32_base_atomics
+
+const sampler_t gInSampler = CLK_NORMALIZED_COORDS_FALSE |
+                             CLK_ADDRESS_CLAMP_TO_EDGE   |
+                             CLK_FILTER_NEAREST;
+
+__kernel void diff(read_only image2d_t baseline, read_only image2d_t test, __global int* result) {
+    int2 coord = (int2)(get_global_id(0), get_global_id(1));
+    uint4 baselinePixel = read_imageui(baseline, gInSampler, coord);
+    uint4 testPixel = read_imageui(test, gInSampler, coord);
+    int4 pixelCompare = baselinePixel == testPixel;
+    if (baselinePixel.x != testPixel.x ||
+        baselinePixel.y != testPixel.y ||
+        baselinePixel.z != testPixel.z ||
+        baselinePixel.w != testPixel.w) {
+
+        atom_inc(result);
+    }
+}

experimental/skpdiff/main.cpp

@@ -0,0 +1,127 @@
+/*
+ * Copyright 2013 Google Inc.
+ *
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#define __NO_STD_VECTOR // Uses cl::vectpr instead of std::vectpr
+#define __NO_STD_STRING // Uses cl::STRING_CLASS instead of std::string
+#include <CL/cl.hpp>
+
+#include "SkOSFile.h"
+#include "SkStream.h"
+#include "SkString.h"
+#include "SkTArray.h"
+#include "SkTDArray.h"
+
+#include "SkImageDiffer.h"
+#include "SkCLImageDiffer.h"
+#include "skpdiff_util.h"
+
+/// A callback for any OpenCL errors
+CL_CALLBACK void error_notify(const char* errorInfo, const void* privateInfoSize, ::size_t cb, void* userData) {
+    SkDebugf("OpenCL error notify: %s\n", errorInfo);
+    exit(1);
+}
+
+/// Creates a device and context with OpenCL
+static bool init_device_and_context(cl::Device* device, cl::Context* context) {
+    // Query for a platform
+    cl::vector<cl::Platform> platformList;
+    cl::Platform::get(&platformList);
+    SkDebugf("The number of platforms is %u\n", platformList.size());
+
+    // Print some information about the platform for debugging
+    cl::Platform& platform = platformList[0];
+    cl::STRING_CLASS platformName;
+    platform.getInfo(CL_PLATFORM_NAME, &platformName);
+    SkDebugf("Platform index 0 is named %s\n", platformName.c_str());
+
+    // Query for a device
+    cl::vector<cl::Device> deviceList;
+    platform.getDevices(CL_DEVICE_TYPE_GPU, &deviceList);
+    SkDebugf("The number of GPU devices is %u\n", deviceList.size());
+
+    // Print some information about the device for debugging
+    *device = deviceList[0];
+    cl::STRING_CLASS deviceName;
+    device->getInfo(CL_DEVICE_NAME, &deviceName);
+    SkDebugf("Device index 0 is named %s\n", deviceName.c_str());
+
+    // Create a CL context and check for all errors
+    cl_int contextErr = CL_SUCCESS;
+    *context = cl::Context(deviceList, NULL, error_notify, NULL, &contextErr);
+    if (contextErr != CL_SUCCESS) {
+        SkDebugf("Context creation failed: %s\n", cl_error_to_string(contextErr));
+        return false;
+    }
+
+    return true;
+}
+
+/// Compares two directories of images with the given differ
+static void diff_directories(const char baselinePath[], const char testPath[], SkImageDiffer* differ) {
+    // Get the files in the baseline, we will then look for those inside the test path
+    SkTArray<SkString> baselineEntries;
+    if (!get_directory(baselinePath, &baselineEntries)) {
+        SkDebugf("Unable to open path \"%s\"\n", baselinePath);
+        return;
+    }
+
+    SkTDArray<int> queuedDiffIDs;
+    for (int baselineIndex = 0; baselineIndex < baselineEntries.count(); baselineIndex++) {
+        const char* baseFilename = baselineEntries[baselineIndex].c_str();
+        SkDebugf("%s\n", baseFilename);
+
+        // Find the real location of each file to compare
+        SkString baselineFile = SkOSPath::SkPathJoin(baselinePath, baseFilename);
+        SkString testFile = SkOSPath::SkPathJoin(testPath, baseFilename);
+
+        // Check that the test file exists and is a file
+        if (sk_exists(testFile.c_str()) && !sk_isdir(testFile.c_str())) {
+            // Queue up the comparison with the differ
+            int diffID = differ->queueDiffOfFile(baselineFile.c_str(), testFile.c_str());
+            if (diffID >= 0) {
+                queuedDiffIDs.push(diffID);
+                SkDebugf("Result: %f\n", differ->getResult(diffID));
+            }
+        } else {
+            SkDebugf("Baseline file \"%s\" has no corresponding test file\n", baselineFile.c_str());
+        }
+    }
+}
+
+static void print_help()
+{
+    SkDebugf(
+    "Usage:\n" \
+    "skpdiff <baseline directory> <test directory>\n\n"
+    );
+}
+
+int main(int argc, char** argv) {
+    if (argc != 3)
+    {
+        print_help();
+        return 1;
+    }
+
+    // Setup OpenCL
+    cl::Device device;
+    cl::Context context;
+    if (!init_device_and_context(&device, &context)) {
+        return 1;
+    }
+
+    // Setup our differ of choice
+    SkCLImageDiffer* differ = SkNEW(SkDifferentPixelsImageDiffer);
+    if (!differ->init(device(), context())) {
+        return 1;
+    }
+
+    // Diff our folders
+    diff_directories(argv[1], argv[2], differ);
+
+    return 0;
+}

experimental/skpdiff/skpdiff.gyp

@@ -0,0 +1,35 @@
+# GYP file to build skpdiff.
+#
+# To build on Linux:
+#  ./gyp_skia skpdiff.gyp && make skpdiff
+#
+{
+  'targets': [
+    {
+      'target_name': 'skpdiff',
+      'type': 'executable',
+      'sources': [
+        'main.cpp',
+        'SkImageDiffer.cpp',
+        'SkCLImageDiffer.cpp',
+        'skpdiff_util.cpp',
+      ],
+      'dependencies': [
+        '../../gyp/skia_lib.gyp:skia_lib',
+      ],
+      'link_settings': {
+        'libraries': [
+          '-lOpenCL',
+        ],
+      },
+    },
+  ],
+  'conditions': [
+  ],
+}
+
+# Local Variables:
+# tab-width:2
+# indent-tabs-mode:nil
+# End:
+# vim: set expandtab tabstop=2 shiftwidth=2:

experimental/skpdiff/skpdiff_util.cpp

@@ -0,0 +1,94 @@
+/*
+ * 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 <time.h>
+#include <dirent.h>
+#include "SkOSFile.h"
+#include "skpdiff_util.h"
+
+const char* cl_error_to_string(cl_int err) {
+    switch (err) {
+        case CL_SUCCESS:                         return "CL_SUCCESS";
+        case CL_DEVICE_NOT_FOUND:                return "CL_DEVICE_NOT_FOUND";
+        case CL_DEVICE_NOT_AVAILABLE:            return "CL_DEVICE_NOT_AVAILABLE";
+        case CL_COMPILER_NOT_AVAILABLE:          return "CL_COMPILER_NOT_AVAILABLE";
+        case CL_MEM_OBJECT_ALLOCATION_FAILURE:   return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
+        case CL_OUT_OF_RESOURCES:                return "CL_OUT_OF_RESOURCES";
+        case CL_OUT_OF_HOST_MEMORY:              return "CL_OUT_OF_HOST_MEMORY";
+        case CL_PROFILING_INFO_NOT_AVAILABLE:    return "CL_PROFILING_INFO_NOT_AVAILABLE";
+        case CL_MEM_COPY_OVERLAP:                return "CL_MEM_COPY_OVERLAP";
+        case CL_IMAGE_FORMAT_MISMATCH:           return "CL_IMAGE_FORMAT_MISMATCH";
+        case CL_IMAGE_FORMAT_NOT_SUPPORTED:      return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
+        case CL_BUILD_PROGRAM_FAILURE:           return "CL_BUILD_PROGRAM_FAILURE";
+        case CL_MAP_FAILURE:                     return "CL_MAP_FAILURE";
+        case CL_INVALID_VALUE:                   return "CL_INVALID_VALUE";
+        case CL_INVALID_DEVICE_TYPE:             return "CL_INVALID_DEVICE_TYPE";
+        case CL_INVALID_PLATFORM:                return "CL_INVALID_PLATFORM";
+        case CL_INVALID_DEVICE:                  return "CL_INVALID_DEVICE";
+        case CL_INVALID_CONTEXT:                 return "CL_INVALID_CONTEXT";
+        case CL_INVALID_QUEUE_PROPERTIES:        return "CL_INVALID_QUEUE_PROPERTIES";
+        case CL_INVALID_COMMAND_QUEUE:           return "CL_INVALID_COMMAND_QUEUE";
+        case CL_INVALID_HOST_PTR:                return "CL_INVALID_HOST_PTR";
+        case CL_INVALID_MEM_OBJECT:              return "CL_INVALID_MEM_OBJECT";
+        case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
+        case CL_INVALID_IMAGE_SIZE:              return "CL_INVALID_IMAGE_SIZE";
+        case CL_INVALID_SAMPLER:                 return "CL_INVALID_SAMPLER";
+        case CL_INVALID_BINARY:                  return "CL_INVALID_BINARY";
+        case CL_INVALID_BUILD_OPTIONS:           return "CL_INVALID_BUILD_OPTIONS";
+        case CL_INVALID_PROGRAM:                 return "CL_INVALID_PROGRAM";
+        case CL_INVALID_PROGRAM_EXECUTABLE:      return "CL_INVALID_PROGRAM_EXECUTABLE";
+        case CL_INVALID_KERNEL_NAME:             return "CL_INVALID_KERNEL_NAME";
+        case CL_INVALID_KERNEL_DEFINITION:       return "CL_INVALID_KERNEL_DEFINITION";
+        case CL_INVALID_KERNEL:                  return "CL_INVALID_KERNEL";
+        case CL_INVALID_ARG_INDEX:               return "CL_INVALID_ARG_INDEX";
+        case CL_INVALID_ARG_VALUE:               return "CL_INVALID_ARG_VALUE";
+        case CL_INVALID_ARG_SIZE:                return "CL_INVALID_ARG_SIZE";
+        case CL_INVALID_KERNEL_ARGS:             return "CL_INVALID_KERNEL_ARGS";
+        case CL_INVALID_WORK_DIMENSION:          return "CL_INVALID_WORK_DIMENSION";
+        case CL_INVALID_WORK_GROUP_SIZE:         return "CL_INVALID_WORK_GROUP_SIZE";
+        case CL_INVALID_WORK_ITEM_SIZE:          return "CL_INVALID_WORK_ITEM_SIZE";
+        case CL_INVALID_GLOBAL_OFFSET:           return "CL_INVALID_GLOBAL_OFFSET";
+        case CL_INVALID_EVENT_WAIT_LIST:         return "CL_INVALID_EVENT_WAIT_LIST";
+        case CL_INVALID_EVENT:                   return "CL_INVALID_EVENT";
+        case CL_INVALID_OPERATION:               return "CL_INVALID_OPERATION";
+        case CL_INVALID_GL_OBJECT:               return "CL_INVALID_GL_OBJECT";
+        case CL_INVALID_BUFFER_SIZE:             return "CL_INVALID_BUFFER_SIZE";
+        case CL_INVALID_MIP_LEVEL:               return "CL_INVALID_MIP_LEVEL";
+        default:                                 return "UNKNOWN";
+    }
+    return "UNKNOWN";
+}
+
+
+double get_seconds() {
+    struct timespec currentTime;
+    clock_gettime(CLOCK_REALTIME, &currentTime);
+    return currentTime.tv_sec + (double)currentTime.tv_nsec / 1e9;
+}
+
+bool get_directory(const char path[], SkTArray<SkString>* entries) {
+    // Open the directory and check for success
+    DIR* dir = opendir(path);
+    if (NULL == dir) {
+        return false;
+    }
+
+    // Loop through dir entries until there are none left (i.e. readdir returns NULL)
+    struct dirent* entry;
+    while ((entry = readdir(dir))) {
+        // dirent only gives relative paths, we need to join them to the base path to check if they
+        // are directories.
+        SkString joinedPath = SkOSPath::SkPathJoin(path, entry->d_name);
+
+        // We only care about files
+        if (!sk_isdir(joinedPath.c_str())) {
+            entries->push_back(SkString(entry->d_name));
+        }
+    }
+
+    return true;
+}

experimental/skpdiff/skpdiff_util.h

@@ -0,0 +1,38 @@
+/*
+ * Copyright 2013 Google Inc.
+ *
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#ifndef skpdiff_util_DEFINED
+#define skpdiff_util_DEFINED
+
+#include <CL/cl.h>
+#include "SkString.h"
+#include "SkTArray.h"
+
+/**
+ * Converts an OpenCL error number into the string of its enumeration name.
+ * @param  err The OpenCL error number
+ * @return The string of the name of the error; "UNKOWN" if the error number is invalid
+ */
+const char* cl_error_to_string(cl_int err);
+
+/**
+ * Get a positive monotonic real-time measure of the amount of seconds since some undefined epoch.
+ * Maximum precision is the goal of this routine.
+ * @return Amount of time in seconds since some epoch
+ */
+double get_seconds();
+
+/**
+ * Get file entries of the given directory.
+ * @param  path    A path to a directory to enumerate
+ * @param  entries A vector to return the results into
+ * @return         True on success, false otherwise
+ */
+bool get_directory(const char path[], SkTArray<SkString>* entries);
+
+
+#endif