1d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com/* 2d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com * Copyright 2013 Google Inc. 3d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com * 4d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com * Use of this source code is governed by a BSD-style license that can be 5d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com * found in the LICENSE file. 6d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com */ 7d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 8d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com#include "SkBitmap.h" 9d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 10d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com#include "SkDifferentPixelsMetric.h" 11d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com#include "skpdiff_util.h" 12d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 13d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.comstatic const char kDifferentPixelsKernelSource[] = 14d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com "#pragma OPENCL_EXTENSION cl_khr_global_int32_base_atomics \n" 15d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " \n" 16d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com "const sampler_t gInSampler = CLK_NORMALIZED_COORDS_FALSE | \n" 17d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " CLK_ADDRESS_CLAMP_TO_EDGE | \n" 18d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " CLK_FILTER_NEAREST; \n" 19d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " \n" 20d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com "__kernel void diff(read_only image2d_t baseline, read_only image2d_t test, \n" 21efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com " __global int* result) { \n" 22d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " int2 coord = (int2)(get_global_id(0), get_global_id(1)); \n" 23d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " uint4 baselinePixel = read_imageui(baseline, gInSampler, coord); \n" 24d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " uint4 testPixel = read_imageui(test, gInSampler, coord); \n" 25d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " if (baselinePixel.x != testPixel.x || \n" 26d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " baselinePixel.y != testPixel.y || \n" 27d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " baselinePixel.z != testPixel.z || \n" 28d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " baselinePixel.w != testPixel.w) { \n" 29d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " \n" 30efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com " atomic_inc(result); \n" 31efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com " // TODO: generate alpha mask \n" 32d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com " } \n" 33d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com "} \n"; 34d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 35efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.comconst char* SkDifferentPixelsMetric::getName() const { 36d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com return "different_pixels"; 37d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com} 38d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 3954f1ad8bb5bdd2ac2ea7981427abeb193383d449epogerbool SkDifferentPixelsMetric::diff(SkBitmap* baseline, SkBitmap* test, bool computeAlphaMask, 4054f1ad8bb5bdd2ac2ea7981427abeb193383d449epoger bool computeRgbDiff, bool computeWhiteDiff, 41efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com Result* result) const { 42d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com double startTime = get_seconds(); 43efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com 44efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com if (!fIsGood) { 45efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com return false; 46efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com } 47d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 48d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // If we never end up running the kernel, include some safe defaults in the result. 49efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com result->poiCount = 0; 50d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 51d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // Ensure the images are comparable 52d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com if (baseline->width() != test->width() || baseline->height() != test->height() || 53d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com baseline->width() <= 0 || baseline->height() <= 0 || 54d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com baseline->config() != test->config()) { 55efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com return false; 56d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com } 57d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 58efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com cl_mem baselineImage; 59efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com cl_mem testImage; 60efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com cl_mem resultsBuffer; 61efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com 62d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // Upload images to the CL device 63efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com if (!this->makeImage2D(baseline, &baselineImage) || !this->makeImage2D(test, &testImage)) { 64efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com SkDebugf("creation of openCL images failed"); 65efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com return false; 66d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com } 67d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 68d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // A small hack that makes calculating percentage difference easier later on. 69efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com result->result = 1.0 / ((double)baseline->width() * baseline->height()); 70d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 71d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // Make a buffer to store results into. It must be initialized with pointers to memory. 72d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com static const int kZero = 0; 73d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // We know OpenCL won't write to it because we use CL_MEM_COPY_HOST_PTR 74efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, 75efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com sizeof(int), (int*)&kZero, NULL); 76d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 77d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // Set all kernel arguments 78efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &baselineImage); 79efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &testImage); 80efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &resultsBuffer); 81d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com if (CL_SUCCESS != setArgErr) { 82d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); 83efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com return false; 84d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com } 85d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 86d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // Queue this diff on the CL device 87d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com cl_event event; 88d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com const size_t workSize[] = { baseline->width(), baseline->height() }; 89d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com cl_int enqueueErr; 90d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSize, 91d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com NULL, 0, NULL, &event); 92d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com if (CL_SUCCESS != enqueueErr) { 93d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); 94efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com return false; 95d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com } 96d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 97d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // This makes things totally synchronous. Actual queue is not ready yet 98d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com clWaitForEvents(1, &event); 99d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 100d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // Immediate read back the results 101efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com clEnqueueReadBuffer(fCommandQueue, resultsBuffer, CL_TRUE, 0, 102efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com sizeof(int), &result->poiCount, 0, NULL, NULL); 103efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com result->result *= (double)result->poiCount; 104efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com result->result = (1.0 - result->result); 105d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 106d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com // Release all the buffers created 107efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com clReleaseMemObject(resultsBuffer); 108efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com clReleaseMemObject(baselineImage); 109efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com clReleaseMemObject(testImage); 110d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 111efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com result->timeElapsed = get_seconds() - startTime; 112efc51b79a22348e3c2596e872609a7a4b018e531djsollen@google.com return true; 113d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com} 114d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 115d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.combool SkDifferentPixelsMetric::onInit() { 116d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel)) { 117d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com return false; 118d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com } 119d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com 120d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com return true; 121d658568a8260b858411e975b1a6011d2a496cc6fzachr@google.com} 122