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