| OLD | NEW |
| (Empty) |
| 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 comp
uteAlphaMask, | |
| 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_HOS
T_PTR, | |
| 75 sizeof(int), (int*)&kZero, nullptr); | |
| 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, nullptr, work
Size, | |
| 91 nullptr, 0, nullptr, &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, nullptr, nullptr); | |
| 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 } | |
| OLD | NEW |