| 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, __global int2* poi) {
\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 " int poiIndex = atomic_inc(result);
\n" | |
| 31 " poi[poiIndex] = coord;
\n" | |
| 32 " }
\n" | |
| 33 "}
\n"; | |
| 34 | |
| 35 struct SkDifferentPixelsMetric::QueuedDiff { | |
| 36 bool finished; | |
| 37 double result; | |
| 38 int numDiffPixels; | |
| 39 SkIPoint* poi; | |
| 40 cl_mem baseline; | |
| 41 cl_mem test; | |
| 42 cl_mem resultsBuffer; | |
| 43 cl_mem poiBuffer; | |
| 44 }; | |
| 45 | |
| 46 const char* SkDifferentPixelsMetric::getName() { | |
| 47 return "different_pixels"; | |
| 48 } | |
| 49 | |
| 50 int SkDifferentPixelsMetric::queueDiff(SkBitmap* baseline, SkBitmap* test) { | |
| 51 int diffID = fQueuedDiffs.count(); | |
| 52 double startTime = get_seconds(); | |
| 53 QueuedDiff* diff = fQueuedDiffs.push(); | |
| 54 | |
| 55 // If we never end up running the kernel, include some safe defaults in the
result. | |
| 56 diff->finished = false; | |
| 57 diff->result = -1.0; | |
| 58 diff->numDiffPixels = 0; | |
| 59 diff->poi = NULL; | |
| 60 | |
| 61 // Ensure the images are comparable | |
| 62 if (baseline->width() != test->width() || baseline->height() != test->height
() || | |
| 63 baseline->width() <= 0 || baseline->height() <= 0 || | |
| 64 baseline->config() != test->config()) { | |
| 65 diff->finished = true; | |
| 66 return diffID; | |
| 67 } | |
| 68 | |
| 69 // Upload images to the CL device | |
| 70 if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test
, &diff->test)) { | |
| 71 diff->finished = true; | |
| 72 fIsGood = false; | |
| 73 return -1; | |
| 74 } | |
| 75 | |
| 76 // A small hack that makes calculating percentage difference easier later on
. | |
| 77 diff->result = 1.0 / ((double)baseline->width() * baseline->height()); | |
| 78 | |
| 79 // Make a buffer to store results into. It must be initialized with pointers
to memory. | |
| 80 static const int kZero = 0; | |
| 81 // We know OpenCL won't write to it because we use CL_MEM_COPY_HOST_PTR | |
| 82 diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_CO
PY_HOST_PTR, | |
| 83 sizeof(int), (int*)&kZero, NULL); | |
| 84 | |
| 85 diff->poiBuffer = clCreateBuffer(fContext, CL_MEM_WRITE_ONLY, | |
| 86 sizeof(int) * 2 * baseline->width() * basel
ine->height(), | |
| 87 NULL, NULL); | |
| 88 | |
| 89 // Set all kernel arguments | |
| 90 cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baselin
e); | |
| 91 setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); | |
| 92 setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->results
Buffer); | |
| 93 setArgErr |= clSetKernelArg(fKernel, 3, sizeof(cl_mem), &diff->poiBuff
er); | |
| 94 if (CL_SUCCESS != setArgErr) { | |
| 95 SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); | |
| 96 fIsGood = false; | |
| 97 return -1; | |
| 98 } | |
| 99 | |
| 100 // Queue this diff on the CL device | |
| 101 cl_event event; | |
| 102 const size_t workSize[] = { baseline->width(), baseline->height() }; | |
| 103 cl_int enqueueErr; | |
| 104 enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSiz
e, | |
| 105 NULL, 0, NULL, &event); | |
| 106 if (CL_SUCCESS != enqueueErr) { | |
| 107 SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); | |
| 108 fIsGood = false; | |
| 109 return -1; | |
| 110 } | |
| 111 | |
| 112 // This makes things totally synchronous. Actual queue is not ready yet | |
| 113 clWaitForEvents(1, &event); | |
| 114 diff->finished = true; | |
| 115 | |
| 116 // Immediate read back the results | |
| 117 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, | |
| 118 sizeof(int), &diff->numDiffPixels, 0, NULL, NULL); | |
| 119 diff->result *= (double)diff->numDiffPixels; | |
| 120 diff->result = (1.0 - diff->result); | |
| 121 | |
| 122 // Reading a buffer of size zero can cause issues on some (Mac) OpenCL platf
orms. | |
| 123 if (diff->numDiffPixels > 0) { | |
| 124 diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels); | |
| 125 clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0, | |
| 126 sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NU
LL, NULL); | |
| 127 } | |
| 128 | |
| 129 // Release all the buffers created | |
| 130 clReleaseMemObject(diff->poiBuffer); | |
| 131 clReleaseMemObject(diff->resultsBuffer); | |
| 132 clReleaseMemObject(diff->baseline); | |
| 133 clReleaseMemObject(diff->test); | |
| 134 | |
| 135 SkDebugf("Time: %f\n", (get_seconds() - startTime)); | |
| 136 | |
| 137 return diffID; | |
| 138 } | |
| 139 | |
| 140 void SkDifferentPixelsMetric::deleteDiff(int id) { | |
| 141 QueuedDiff* diff = &fQueuedDiffs[id]; | |
| 142 if (NULL != diff->poi) { | |
| 143 SkDELETE_ARRAY(diff->poi); | |
| 144 diff->poi = NULL; | |
| 145 } | |
| 146 } | |
| 147 | |
| 148 bool SkDifferentPixelsMetric::isFinished(int id) { | |
| 149 return fQueuedDiffs[id].finished; | |
| 150 } | |
| 151 | |
| 152 double SkDifferentPixelsMetric::getResult(int id) { | |
| 153 return fQueuedDiffs[id].result; | |
| 154 } | |
| 155 | |
| 156 int SkDifferentPixelsMetric::getPointsOfInterestCount(int id) { | |
| 157 return fQueuedDiffs[id].numDiffPixels; | |
| 158 } | |
| 159 | |
| 160 SkIPoint* SkDifferentPixelsMetric::getPointsOfInterest(int id) { | |
| 161 return fQueuedDiffs[id].poi; | |
| 162 } | |
| 163 | |
| 164 bool SkDifferentPixelsMetric::onInit() { | |
| 165 if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel))
{ | |
| 166 return false; | |
| 167 } | |
| 168 | |
| 169 return true; | |
| 170 } | |
| OLD | NEW |
|---|
Chromium Code Reviews
