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