| OLD | NEW |
| 1 /* | 1 /* |
| 2 * Copyright 2013 Google Inc. | 2 * Copyright 2013 Google Inc. |
| 3 * | 3 * |
| 4 * Use of this source code is governed by a BSD-style license that can be | 4 * Use of this source code is governed by a BSD-style license that can be |
| 5 * found in the LICENSE file. | 5 * found in the LICENSE file. |
| 6 */ | 6 */ |
| 7 | 7 |
| 8 #include "SkBitmap.h" | 8 #include "SkBitmap.h" |
| 9 | 9 |
| 10 #include "SkDifferentPixelsMetric.h" | 10 #include "SkDifferentPixelsMetric.h" |
| 11 #include "skpdiff_util.h" | 11 #include "skpdiff_util.h" |
| 12 | 12 |
| 13 static const char kDifferentPixelsKernelSource[] = | 13 static const char kDifferentPixelsKernelSource[] = |
| 14 "#pragma OPENCL_EXTENSION cl_khr_global_int32_base_atomics
\n" | 14 "#pragma OPENCL_EXTENSION cl_khr_global_int32_base_atomics
\n" |
| 15 "
\n" | 15 "
\n" |
| 16 "const sampler_t gInSampler = CLK_NORMALIZED_COORDS_FALSE |
\n" | 16 "const sampler_t gInSampler = CLK_NORMALIZED_COORDS_FALSE |
\n" |
| 17 " CLK_ADDRESS_CLAMP_TO_EDGE |
\n" | 17 " CLK_ADDRESS_CLAMP_TO_EDGE |
\n" |
| 18 " CLK_FILTER_NEAREST;
\n" | 18 " CLK_FILTER_NEAREST;
\n" |
| 19 "
\n" | 19 "
\n" |
| 20 "__kernel void diff(read_only image2d_t baseline, read_only image2d_t test,
\n" | 20 "__kernel void diff(read_only image2d_t baseline, read_only image2d_t test,
\n" |
| 21 " __global int* result, __global int2* poi) {
\n" | 21 " __global int* result, __global int2* poi) {
\n" |
| 22 " int2 coord = (int2)(get_global_id(0), get_global_id(1));
\n" | 22 " int2 coord = (int2)(get_global_id(0), get_global_id(1));
\n" |
| 23 " uint4 baselinePixel = read_imageui(baseline, gInSampler, coord);
\n" | 23 " uint4 baselinePixel = read_imageui(baseline, gInSampler, coord);
\n" |
| 24 " uint4 testPixel = read_imageui(test, gInSampler, coord);
\n" | 24 " uint4 testPixel = read_imageui(test, gInSampler, coord);
\n" |
| 25 " int4 pixelCompare = baselinePixel == testPixel;
\n" | |
| 26 " if (baselinePixel.x != testPixel.x ||
\n" | 25 " if (baselinePixel.x != testPixel.x ||
\n" |
| 27 " baselinePixel.y != testPixel.y ||
\n" | 26 " baselinePixel.y != testPixel.y ||
\n" |
| 28 " baselinePixel.z != testPixel.z ||
\n" | 27 " baselinePixel.z != testPixel.z ||
\n" |
| 29 " baselinePixel.w != testPixel.w) {
\n" | 28 " baselinePixel.w != testPixel.w) {
\n" |
| 30 "
\n" | 29 "
\n" |
| 31 " int poiIndex = atomic_inc(result);
\n" | 30 " int poiIndex = atomic_inc(result);
\n" |
| 32 " poi[poiIndex] = coord;
\n" | 31 " poi[poiIndex] = coord;
\n" |
| 33 " }
\n" | 32 " }
\n" |
| 34 "}
\n"; | 33 "}
\n"; |
| 35 | 34 |
| (...skipping 77 matching lines...) Expand 10 before | Expand all | Expand 10 after Loading... |
| 113 // This makes things totally synchronous. Actual queue is not ready yet | 112 // This makes things totally synchronous. Actual queue is not ready yet |
| 114 clWaitForEvents(1, &event); | 113 clWaitForEvents(1, &event); |
| 115 diff->finished = true; | 114 diff->finished = true; |
| 116 | 115 |
| 117 // Immediate read back the results | 116 // Immediate read back the results |
| 118 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, | 117 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, |
| 119 sizeof(int), &diff->numDiffPixels, 0, NULL, NULL); | 118 sizeof(int), &diff->numDiffPixels, 0, NULL, NULL); |
| 120 diff->result *= (double)diff->numDiffPixels; | 119 diff->result *= (double)diff->numDiffPixels; |
| 121 diff->result = (1.0 - diff->result); | 120 diff->result = (1.0 - diff->result); |
| 122 | 121 |
| 123 diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels); | 122 // Reading a buffer of size zero can cause issues on some (Mac) OpenCL platf
orms. |
| 124 clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0, | 123 if (diff->numDiffPixels > 0) { |
| 124 diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels); |
| 125 clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0, |
| 125 sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NU
LL, NULL); | 126 sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NU
LL, NULL); |
| 127 } |
| 126 | 128 |
| 127 // Release all the buffers created | 129 // Release all the buffers created |
| 128 clReleaseMemObject(diff->poiBuffer); | 130 clReleaseMemObject(diff->poiBuffer); |
| 129 clReleaseMemObject(diff->resultsBuffer); | 131 clReleaseMemObject(diff->resultsBuffer); |
| 130 clReleaseMemObject(diff->baseline); | 132 clReleaseMemObject(diff->baseline); |
| 131 clReleaseMemObject(diff->test); | 133 clReleaseMemObject(diff->test); |
| 132 | 134 |
| 133 SkDebugf("Time: %f\n", (get_seconds() - startTime)); | 135 SkDebugf("Time: %f\n", (get_seconds() - startTime)); |
| 134 | 136 |
| 135 return diffID; | 137 return diffID; |
| (...skipping 23 matching lines...) Expand all Loading... |
| 159 return fQueuedDiffs[id].poi; | 161 return fQueuedDiffs[id].poi; |
| 160 } | 162 } |
| 161 | 163 |
| 162 bool SkDifferentPixelsMetric::onInit() { | 164 bool SkDifferentPixelsMetric::onInit() { |
| 163 if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel))
{ | 165 if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel))
{ |
| 164 return false; | 166 return false; |
| 165 } | 167 } |
| 166 | 168 |
| 167 return true; | 169 return true; |
| 168 } | 170 } |
| OLD | NEW |