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