src/third_party/skia/tools/skpdiff/SkDifferentPixelsMetric_opencl.cpp

   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) {                                 \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     "        atomic_inc(result);                                                \n"
  31     "        // TODO: generate alpha mask                                       \n"
  32     "    }                                                                      \n"
  33     "}                                                                          \n";
  34
  35 const char* SkDifferentPixelsMetric::getName() const {
  36     return "different_pixels";
  37 }
  38
  39 bool SkDifferentPixelsMetric::diff(SkBitmap* baseline, SkBitmap* test, bool computeAlphaMask,
  40                                    bool computeRgbDiff, bool computeWhiteDiff,
  41                                    Result* result) const {
  42     double startTime = get_seconds();
  43
  44     if (!fIsGood) {
  45         return false;
  46     }
  47
  48     // If we never end up running the kernel, include some safe defaults in the result.
  49     result->poiCount = 0;
  50
  51     // Ensure the images are comparable
  52     if (baseline->width() != test->width() || baseline->height() != test->height() ||
  53         baseline->width() <= 0 || baseline->height() <= 0 ||
  54         baseline->config() != test->config()) {
  55         return false;
  56     }
  57
  58     cl_mem baselineImage;
  59     cl_mem testImage;
  60     cl_mem resultsBuffer;
  61
  62     // Upload images to the CL device
  63     if (!this->makeImage2D(baseline, &baselineImage) || !this->makeImage2D(test, &testImage)) {
  64         SkDebugf("creation of openCL images failed");
  65         return false;
  66     }
  67
  68     // A small hack that makes calculating percentage difference easier later on.
  69     result->result = 1.0 / ((double)baseline->width() * baseline->height());
  70
  71     // Make a buffer to store results into. It must be initialized with pointers to memory.
  72     static const int kZero = 0;
  73     // We know OpenCL won't write to it because we use CL_MEM_COPY_HOST_PTR
  74     resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
  75                                    sizeof(int), (int*)&kZero, NULL);
  76
  77     // Set all kernel arguments
  78     cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &baselineImage);
  79     setArgErr       |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &testImage);
  80     setArgErr       |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &resultsBuffer);
  81     if (CL_SUCCESS != setArgErr) {
  82         SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr));
  83         return false;
  84     }
  85
  86     // Queue this diff on the CL device
  87     cl_event event;
  88     const size_t workSize[] = { baseline->width(), baseline->height() };
  89     cl_int enqueueErr;
  90     enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSize,
  91                                         NULL, 0, NULL, &event);
  92     if (CL_SUCCESS != enqueueErr) {
  93         SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr));
  94         return false;
  95     }
  96
  97     // This makes things totally synchronous. Actual queue is not ready yet
  98     clWaitForEvents(1, &event);
  99
 100     // Immediate read back the results
 101     clEnqueueReadBuffer(fCommandQueue, resultsBuffer, CL_TRUE, 0,
 102                         sizeof(int), &result->poiCount, 0, NULL, NULL);
 103     result->result *= (double)result->poiCount;
 104     result->result = (1.0 - result->result);
 105
 106     // Release all the buffers created
 107     clReleaseMemObject(resultsBuffer);
 108     clReleaseMemObject(baselineImage);
 109     clReleaseMemObject(testImage);
 110
 111     result->timeElapsed = get_seconds() - startTime;
 112     return true;
 113 }
 114
 115 bool SkDifferentPixelsMetric::onInit() {
 116     if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel)) {
 117         return false;
 118     }
 119
 120     return true;
 121 }