fix multithread related crashes in skpdiff

tools/skpdiff/SkDifferentPixelsMetric_opencl.cpp

 /*
  * 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[] =

[mtklein, 2013/11/08 15:06:13]

Out of curiosity, do you see us realistically using the opencl version of this? 
Seems neat, but might be easier to just blow it away if we're only really going
to use the _cpu path.

[djsollen, 2013/11/12 16:40:43]

I don't want to blow it away yet, until I can do some due diligence and know
that the perf benefit is not worth the code complexity.

     "#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"
+    "                   __global int* result) {                                 \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"
     "    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"
+    "        atomic_inc(result);                                                \n"
+    "        // TODO: generate alpha mask                                       \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";
 }
 
 bool SkDifferentPixelsMetric::enablePOIAlphaMask() {
     return false;
 }
 
-int SkDifferentPixelsMetric::queueDiff(SkBitmap* baseline, SkBitmap* test) {
-    int diffID = fQueuedDiffs.count();
+bool SkDifferentPixelsMetric::diff(SkBitmap* baseline, SkBitmap* test, Result* result) {
     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;
+    result->poiCount = 0;
 
     // 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;
+        return false;
     }
 
+    cl_mem baselineImage;
+    cl_mem testImage;
+    cl_mem resultsBuffer;
+
     // Upload images to the CL device
-    if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test, &diff->test)) {
-        diff->finished = true;
-        fIsGood = false;
-        return -1;
+    if (!this->makeImage2D(baseline, &baselineImage) || !this->makeImage2D(test, &testImage)) {
+        SkDebugf("creation of openCL images failed");
+        return false;
     }
 
     // A small hack that makes calculating percentage difference easier later on.
-    diff->result = 1.0 / ((double)baseline->width() * baseline->height());
+    result->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);
+    resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
+                                   sizeof(int), (int*)&kZero, 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);
+    cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &baselineImage);
+    setArgErr       |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &testImage);
+    setArgErr       |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &resultsBuffer);
     if (CL_SUCCESS != setArgErr) {
         SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr));
         fIsGood = false;
-        return -1;
+        return false;
     }
 
     // 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;
+        return false;
     }
 
     // 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);
-
-    // Reading a buffer of size zero can cause issues on some (Mac) OpenCL platforms.
-    if (diff->numDiffPixels > 0) {
-        diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels);
-        clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0,
-                        sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NULL, NULL);
-    }
+    clEnqueueReadBuffer(fCommandQueue, resultsBuffer, CL_TRUE, 0,
+                        sizeof(int), &result->poiCount, 0, NULL, NULL);
+    result->result *= (double)result->poiCount;
+    result->result = (1.0 - result->result);
 
     // Release all the buffers created
-    clReleaseMemObject(diff->poiBuffer);
-    clReleaseMemObject(diff->resultsBuffer);
-    clReleaseMemObject(diff->baseline);
-    clReleaseMemObject(diff->test);
+    clReleaseMemObject(resultsBuffer);
+    clReleaseMemObject(baselineImage);
+    clReleaseMemObject(testImage);
 
-    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;
+    result->timeElapsed = get_seconds() - startTime;
+    return true;
 }
 
 bool SkDifferentPixelsMetric::onInit() {
     if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel)) {
         return false;
     }
 
     return true;
 }