Home | History | Annotate | Download | only in skpdiff 
      1 /*
      2  * Copyright 2013 Google Inc.
      3  *
      4  * Use of this source code is governed by a BSD-style license that can be
      5  * found in the LICENSE file.
      6  */
      7 
      8 #include "SkBitmap.h"
      9 
     10 #include "SkDifferentPixelsMetric.h"
     11 #include "skpdiff_util.h"
     12 
     13 static const char kDifferentPixelsKernelSource[] =
     14     "#pragma OPENCL_EXTENSION cl_khr_global_int32_base_atomics                  \n"
     15     "                                                                           \n"
     16     "const sampler_t gInSampler = CLK_NORMALIZED_COORDS_FALSE |                 \n"
     17     "                             CLK_ADDRESS_CLAMP_TO_EDGE   |                 \n"
     18     "                             CLK_FILTER_NEAREST;                           \n"
     19     "                                                                           \n"
     20     "__kernel void diff(read_only image2d_t baseline, read_only image2d_t test, \n"
     21     "                   __global int* result, __global int2* poi) {             \n"
     22     "    int2 coord = (int2)(get_global_id(0), get_global_id(1));               \n"
     23     "    uint4 baselinePixel = read_imageui(baseline, gInSampler, coord);       \n"
     24     "    uint4 testPixel = read_imageui(test, gInSampler, coord);               \n"
     25     "    if (baselinePixel.x != testPixel.x ||                                  \n"
     26     "        baselinePixel.y != testPixel.y ||                                  \n"
     27     "        baselinePixel.z != testPixel.z ||                                  \n"
     28     "        baselinePixel.w != testPixel.w) {                                  \n"
     29     "                                                                           \n"
     30     "        int poiIndex = atomic_inc(result);                                 \n"
     31     "        poi[poiIndex] = coord;                                             \n"
     32     "    }                                                                      \n"
     33     "}                                                                          \n";
     34 
     35 struct SkDifferentPixelsMetric::QueuedDiff {
     36     bool finished;
     37     double result;
     38     int numDiffPixels;
     39     SkIPoint* poi;
     40     cl_mem baseline;
     41     cl_mem test;
     42     cl_mem resultsBuffer;
     43     cl_mem poiBuffer;
     44 };
     45 
     46 const char* SkDifferentPixelsMetric::getName() {
     47     return "different_pixels";
     48 }
     49 
     50 int SkDifferentPixelsMetric::queueDiff(SkBitmap* baseline, SkBitmap* test) {
     51     int diffID = fQueuedDiffs.count();
     52     double startTime = get_seconds();
     53     QueuedDiff* diff = fQueuedDiffs.push();
     54 
     55     // If we never end up running the kernel, include some safe defaults in the result.
     56     diff->finished = false;
     57     diff->result = -1.0;
     58     diff->numDiffPixels = 0;
     59     diff->poi = NULL;
     60 
     61     // Ensure the images are comparable
     62     if (baseline->width() != test->width() || baseline->height() != test->height() ||
     63         baseline->width() <= 0 || baseline->height() <= 0 ||
     64         baseline->config() != test->config()) {
     65         diff->finished = true;
     66         return diffID;
     67     }
     68 
     69     // Upload images to the CL device
     70     if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test, &diff->test)) {
     71         diff->finished = true;
     72         fIsGood = false;
     73         return -1;
     74     }
     75 
     76     // A small hack that makes calculating percentage difference easier later on.
     77     diff->result = 1.0 / ((double)baseline->width() * baseline->height());
     78 
     79     // Make a buffer to store results into. It must be initialized with pointers to memory.
     80     static const int kZero = 0;
     81     // We know OpenCL won't write to it because we use CL_MEM_COPY_HOST_PTR
     82     diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
     83                                          sizeof(int), (int*)&kZero, NULL);
     84 
     85     diff->poiBuffer = clCreateBuffer(fContext, CL_MEM_WRITE_ONLY,
     86                                      sizeof(int) * 2 * baseline->width() * baseline->height(),
     87                                      NULL, NULL);
     88 
     89     // Set all kernel arguments
     90     cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baseline);
     91     setArgErr       |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test);
     92     setArgErr       |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->resultsBuffer);
     93     setArgErr       |= clSetKernelArg(fKernel, 3, sizeof(cl_mem), &diff->poiBuffer);
     94     if (CL_SUCCESS != setArgErr) {
     95         SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr));
     96         fIsGood = false;
     97         return -1;
     98     }
     99 
    100     // Queue this diff on the CL device
    101     cl_event event;
    102     const size_t workSize[] = { baseline->width(), baseline->height() };
    103     cl_int enqueueErr;
    104     enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSize,
    105                                         NULL, 0, NULL, &event);
    106     if (CL_SUCCESS != enqueueErr) {
    107         SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr));
    108         fIsGood = false;
    109         return -1;
    110     }
    111 
    112     // This makes things totally synchronous. Actual queue is not ready yet
    113     clWaitForEvents(1, &event);
    114     diff->finished = true;
    115 
    116     // Immediate read back the results
    117     clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0,
    118                         sizeof(int), &diff->numDiffPixels, 0, NULL, NULL);
    119     diff->result *= (double)diff->numDiffPixels;
    120     diff->result = (1.0 - diff->result);
    121 
    122     // Reading a buffer of size zero can cause issues on some (Mac) OpenCL platforms.
    123     if (diff->numDiffPixels > 0) {
    124         diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels);
    125         clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0,
    126                         sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NULL, NULL);
    127     }
    128 
    129     // Release all the buffers created
    130     clReleaseMemObject(diff->poiBuffer);
    131     clReleaseMemObject(diff->resultsBuffer);
    132     clReleaseMemObject(diff->baseline);
    133     clReleaseMemObject(diff->test);
    134 
    135     SkDebugf("Time: %f\n", (get_seconds() - startTime));
    136 
    137     return diffID;
    138 }
    139 
    140 void SkDifferentPixelsMetric::deleteDiff(int id) {
    141     QueuedDiff* diff = &fQueuedDiffs[id];
    142     if (NULL != diff->poi) {
    143         SkDELETE_ARRAY(diff->poi);
    144         diff->poi = NULL;
    145     }
    146 }
    147 
    148 bool SkDifferentPixelsMetric::isFinished(int id) {
    149     return fQueuedDiffs[id].finished;
    150 }
    151 
    152 double SkDifferentPixelsMetric::getResult(int id) {
    153     return fQueuedDiffs[id].result;
    154 }
    155 
    156 int SkDifferentPixelsMetric::getPointsOfInterestCount(int id) {
    157     return fQueuedDiffs[id].numDiffPixels;
    158 }
    159 
    160 SkIPoint* SkDifferentPixelsMetric::getPointsOfInterest(int id) {
    161     return fQueuedDiffs[id].poi;
    162 }
    163 
    164 bool SkDifferentPixelsMetric::onInit() {
    165     if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel)) {
    166         return false;
    167     }
    168 
    169     return true;
    170 }
    171