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 |