OLD | NEW |
1 | |
2 /* | 1 /* |
3 * Copyright 2013 Google Inc. | 2 * Copyright 2013 Google Inc. |
4 * | 3 * |
5 * 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 |
6 * found in the LICENSE file. | 5 * found in the LICENSE file. |
7 */ | 6 */ |
8 | 7 |
9 #include <cstring> | 8 #include <cstring> |
10 | 9 |
11 #include "SkBitmap.h" | 10 #include "SkBitmap.h" |
12 #include "SkStream.h" | 11 #include "SkStream.h" |
13 | 12 |
14 #include "SkCLImageDiffer.h" | 13 #include "SkCLImageDiffer.h" |
15 #include "skpdiff_util.h" | 14 #include "skpdiff_util.h" |
16 | 15 |
17 SkCLImageDiffer::SkCLImageDiffer() { | 16 SkCLImageDiffer::SkCLImageDiffer() { |
18 fIsGood = false; | 17 fIsGood = false; |
19 } | 18 } |
20 | 19 |
21 | |
22 bool SkCLImageDiffer::init(cl_device_id device, cl_context context) { | 20 bool SkCLImageDiffer::init(cl_device_id device, cl_context context) { |
23 fContext = context; | 21 fContext = context; |
24 fDevice = device; | 22 fDevice = device; |
25 | 23 |
26 cl_int queueErr; | 24 cl_int queueErr; |
27 fCommandQueue = clCreateCommandQueue(fContext, fDevice, 0, &queueErr); | 25 fCommandQueue = clCreateCommandQueue(fContext, fDevice, 0, &queueErr); |
28 if (CL_SUCCESS != queueErr) { | 26 if (CL_SUCCESS != queueErr) { |
29 SkDebugf("Command queue creation failed: %s\n", cl_error_to_string(queue
Err)); | 27 SkDebugf("Command queue creation failed: %s\n", cl_error_to_string(queue
Err)); |
30 fIsGood = false; | 28 fIsGood = false; |
31 return false; | 29 return false; |
(...skipping 83 matching lines...) Expand 10 before | Expand all | Expand 10 after Loading... |
115 &imageErr); | 113 &imageErr); |
116 bitmap->unlockPixels(); | 114 bitmap->unlockPixels(); |
117 | 115 |
118 if (CL_SUCCESS != imageErr) { | 116 if (CL_SUCCESS != imageErr) { |
119 SkDebugf("Input image creation failed: %s\n", cl_error_to_string(imageEr
r)); | 117 SkDebugf("Input image creation failed: %s\n", cl_error_to_string(imageEr
r)); |
120 return false; | 118 return false; |
121 } | 119 } |
122 | 120 |
123 return true; | 121 return true; |
124 } | 122 } |
125 | |
126 | |
127 //////////////////////////////////////////////////////////////// | |
128 | |
129 struct SkDifferentPixelsImageDiffer::QueuedDiff { | |
130 bool finished; | |
131 double result; | |
132 int numDiffPixels; | |
133 SkIPoint* poi; | |
134 cl_mem baseline; | |
135 cl_mem test; | |
136 cl_mem resultsBuffer; | |
137 cl_mem poiBuffer; | |
138 }; | |
139 | |
140 const char* SkDifferentPixelsImageDiffer::getName() { | |
141 return "different_pixels"; | |
142 } | |
143 | |
144 int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test
) { | |
145 int diffID = fQueuedDiffs.count(); | |
146 double startTime = get_seconds(); | |
147 QueuedDiff* diff = fQueuedDiffs.push(); | |
148 | |
149 // If we never end up running the kernel, include some safe defaults in the
result. | |
150 diff->finished = false; | |
151 diff->result = -1.0; | |
152 diff->numDiffPixels = 0; | |
153 diff->poi = NULL; | |
154 | |
155 // Ensure the images are comparable | |
156 if (baseline->width() != test->width() || baseline->height() != test->height
() || | |
157 baseline->width() <= 0 || baseline->height() <= 0) { | |
158 diff->finished = true; | |
159 return diffID; | |
160 } | |
161 | |
162 // Upload images to the CL device | |
163 if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test
, &diff->test)) { | |
164 diff->finished = true; | |
165 fIsGood = false; | |
166 return -1; | |
167 } | |
168 | |
169 // A small hack that makes calculating percentage difference easier later on
. | |
170 diff->result = 1.0 / ((double)baseline->width() * baseline->height()); | |
171 | |
172 // Make a buffer to store results into. It must be initialized with pointers
to memory. | |
173 static const int kZero = 0; | |
174 // We know OpenCL won't write to it because we use CL_MEM_COPY_HOST_PTR | |
175 diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_CO
PY_HOST_PTR, | |
176 sizeof(int), (int*)&kZero, NULL); | |
177 | |
178 diff->poiBuffer = clCreateBuffer(fContext, CL_MEM_WRITE_ONLY, | |
179 sizeof(int) * 2 * baseline->width() * basel
ine->height(), | |
180 NULL, NULL); | |
181 | |
182 // Set all kernel arguments | |
183 cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baselin
e); | |
184 setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); | |
185 setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->results
Buffer); | |
186 setArgErr |= clSetKernelArg(fKernel, 3, sizeof(cl_mem), &diff->poiBuff
er); | |
187 if (CL_SUCCESS != setArgErr) { | |
188 SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); | |
189 fIsGood = false; | |
190 return -1; | |
191 } | |
192 | |
193 // Queue this diff on the CL device | |
194 cl_event event; | |
195 const size_t workSize[] = { baseline->width(), baseline->height() }; | |
196 cl_int enqueueErr; | |
197 enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSiz
e, | |
198 NULL, 0, NULL, &event); | |
199 if (CL_SUCCESS != enqueueErr) { | |
200 SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); | |
201 fIsGood = false; | |
202 return -1; | |
203 } | |
204 | |
205 // This makes things totally synchronous. Actual queue is not ready yet | |
206 clWaitForEvents(1, &event); | |
207 diff->finished = true; | |
208 | |
209 // Immediate read back the results | |
210 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, | |
211 sizeof(int), &diff->numDiffPixels, 0, NULL, NULL); | |
212 diff->result *= (double)diff->numDiffPixels; | |
213 diff->result = (1.0 - diff->result); | |
214 | |
215 diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels); | |
216 clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0, | |
217 sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NU
LL, NULL); | |
218 | |
219 // Release all the buffers created | |
220 clReleaseMemObject(diff->poiBuffer); | |
221 clReleaseMemObject(diff->resultsBuffer); | |
222 clReleaseMemObject(diff->baseline); | |
223 clReleaseMemObject(diff->test); | |
224 | |
225 SkDebugf("Time: %f\n", (get_seconds() - startTime)); | |
226 | |
227 return diffID; | |
228 } | |
229 | |
230 void SkDifferentPixelsImageDiffer::deleteDiff(int id) { | |
231 QueuedDiff* diff = &fQueuedDiffs[id]; | |
232 if (NULL != diff->poi) { | |
233 SkDELETE_ARRAY(diff->poi); | |
234 diff->poi = NULL; | |
235 } | |
236 } | |
237 | |
238 bool SkDifferentPixelsImageDiffer::isFinished(int id) { | |
239 return fQueuedDiffs[id].finished; | |
240 } | |
241 | |
242 double SkDifferentPixelsImageDiffer::getResult(int id) { | |
243 return fQueuedDiffs[id].result; | |
244 } | |
245 | |
246 int SkDifferentPixelsImageDiffer::getPointsOfInterestCount(int id) { | |
247 return fQueuedDiffs[id].numDiffPixels; | |
248 } | |
249 | |
250 SkIPoint* SkDifferentPixelsImageDiffer::getPointsOfInterest(int id) { | |
251 return fQueuedDiffs[id].poi; | |
252 } | |
253 | |
254 bool SkDifferentPixelsImageDiffer::onInit() { | |
255 if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel)
) { | |
256 return false; | |
257 } | |
258 | |
259 return true; | |
260 } | |
OLD | NEW |