OLD | NEW |
1 | 1 |
2 /* | 2 /* |
3 * Copyright 2013 Google Inc. | 3 * Copyright 2013 Google Inc. |
4 * | 4 * |
5 * Use of this source code is governed by a BSD-style license that can be | 5 * Use of this source code is governed by a BSD-style license that can be |
6 * found in the LICENSE file. | 6 * found in the LICENSE file. |
7 */ | 7 */ |
8 | 8 |
9 #include <cstring> | 9 #include <cstring> |
10 | 10 |
(...skipping 51 matching lines...) Expand 10 before | Expand all | Expand 10 after Loading... |
62 bool SkCLImageDiffer::loadKernelSource(const char source[], const char name[], c
l_kernel* kernel) { | 62 bool SkCLImageDiffer::loadKernelSource(const char source[], const char name[], c
l_kernel* kernel) { |
63 // Build the kernel source | 63 // Build the kernel source |
64 size_t sourceLen = strlen(source); | 64 size_t sourceLen = strlen(source); |
65 cl_program program = clCreateProgramWithSource(fContext, 1, &source, &source
Len, NULL); | 65 cl_program program = clCreateProgramWithSource(fContext, 1, &source, &source
Len, NULL); |
66 cl_int programErr = clBuildProgram(program, 1, &fDevice, "", NULL, NULL); | 66 cl_int programErr = clBuildProgram(program, 1, &fDevice, "", NULL, NULL); |
67 if (CL_SUCCESS != programErr) { | 67 if (CL_SUCCESS != programErr) { |
68 SkDebugf("Program creation failed: %s\n", cl_error_to_string(programErr)
); | 68 SkDebugf("Program creation failed: %s\n", cl_error_to_string(programErr)
); |
69 | 69 |
70 // Attempt to get information about why the build failed | 70 // Attempt to get information about why the build failed |
71 char buildLog[4096]; | 71 char buildLog[4096]; |
72 clGetProgramBuildInfo(program, fDevice, CL_PROGRAM_BUILD_LOG, sizeof(bui
ldLog), buildLog, NULL); | 72 clGetProgramBuildInfo(program, fDevice, CL_PROGRAM_BUILD_LOG, sizeof(bui
ldLog), |
| 73 buildLog, NULL); |
73 SkDebugf("Build log: %s\n", buildLog); | 74 SkDebugf("Build log: %s\n", buildLog); |
74 | 75 |
75 return false; | 76 return false; |
76 } | 77 } |
77 | 78 |
78 cl_int kernelErr; | 79 cl_int kernelErr; |
79 *kernel = clCreateKernel(program, name, &kernelErr); | 80 *kernel = clCreateKernel(program, name, &kernelErr); |
80 if (CL_SUCCESS != kernelErr) { | 81 if (CL_SUCCESS != kernelErr) { |
81 SkDebugf("Kernel creation failed: %s\n", cl_error_to_string(kernelErr)); | 82 SkDebugf("Kernel creation failed: %s\n", cl_error_to_string(kernelErr)); |
82 return false; | 83 return false; |
(...skipping 35 matching lines...) Expand 10 before | Expand all | Expand 10 after Loading... |
118 SkDebugf("Input image creation failed: %s\n", cl_error_to_string(imageEr
r)); | 119 SkDebugf("Input image creation failed: %s\n", cl_error_to_string(imageEr
r)); |
119 return false; | 120 return false; |
120 } | 121 } |
121 | 122 |
122 return true; | 123 return true; |
123 } | 124 } |
124 | 125 |
125 | 126 |
126 //////////////////////////////////////////////////////////////// | 127 //////////////////////////////////////////////////////////////// |
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 |
128 const char* SkDifferentPixelsImageDiffer::getName() { | 140 const char* SkDifferentPixelsImageDiffer::getName() { |
129 return "different_pixels"; | 141 return "different_pixels"; |
130 } | 142 } |
131 | 143 |
132 int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test
) { | 144 int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test
) { |
133 int diffID = fQueuedDiffs.count(); | 145 int diffID = fQueuedDiffs.count(); |
134 double startTime = get_seconds(); | 146 double startTime = get_seconds(); |
135 QueuedDiff* diff = fQueuedDiffs.push(); | 147 QueuedDiff* diff = fQueuedDiffs.push(); |
136 | 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 |
137 // Ensure the images are comparable | 155 // Ensure the images are comparable |
138 if (baseline->width() != test->width() || baseline->height() != test->height
() || | 156 if (baseline->width() != test->width() || baseline->height() != test->height
() || |
139 baseline->width() <= 0 || baseline->height() <= 0) { | 157 baseline->width() <= 0 || baseline->height() <= 0) { |
140 diff->finished = true; | 158 diff->finished = true; |
141 diff->result = 0.0; | |
142 return diffID; | 159 return diffID; |
143 } | 160 } |
144 | 161 |
145 // Upload images to the CL device | 162 // Upload images to the CL device |
146 if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test
, &diff->test)) { | 163 if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test
, &diff->test)) { |
147 diff->finished = true; | 164 diff->finished = true; |
148 diff->result = 0.0; | |
149 fIsGood = false; | 165 fIsGood = false; |
150 return -1; | 166 return -1; |
151 } | 167 } |
152 | 168 |
153 // A small hack that makes calculating percentage difference easier later on
. | 169 // A small hack that makes calculating percentage difference easier later on
. |
154 diff->result = 1.0 / ((double)baseline->width() * baseline->height()); | 170 diff->result = 1.0 / ((double)baseline->width() * baseline->height()); |
155 | 171 |
156 // Make a buffer to store results into | 172 // Make a buffer to store results into. It must be initialized with pointers
to memory. |
157 int numDiffPixels = 0; | 173 static const int kZero = 0; |
| 174 // We know OpenCL won't write to it because we use CL_MEM_COPY_HOST_PTR |
158 diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_CO
PY_HOST_PTR, | 175 diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_CO
PY_HOST_PTR, |
159 sizeof(int), &numDiffPixels, NULL); | 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); |
160 | 181 |
161 // Set all kernel arguments | 182 // Set all kernel arguments |
162 cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baselin
e); | 183 cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baselin
e); |
163 setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); | 184 setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); |
164 setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->results
Buffer); | 185 setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->results
Buffer); |
| 186 setArgErr |= clSetKernelArg(fKernel, 3, sizeof(cl_mem), &diff->poiBuff
er); |
165 if (CL_SUCCESS != setArgErr) { | 187 if (CL_SUCCESS != setArgErr) { |
166 SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); | 188 SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); |
167 fIsGood = false; | 189 fIsGood = false; |
168 return -1; | 190 return -1; |
169 } | 191 } |
170 | 192 |
171 // Queue this diff on the CL device | 193 // Queue this diff on the CL device |
172 cl_event event; | 194 cl_event event; |
173 const size_t workSize[] = { baseline->width(), baseline->height() }; | 195 const size_t workSize[] = { baseline->width(), baseline->height() }; |
174 cl_int enqueueErr; | 196 cl_int enqueueErr; |
175 enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSiz
e, NULL, 0, NULL, &event); | 197 enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSiz
e, |
| 198 NULL, 0, NULL, &event); |
176 if (CL_SUCCESS != enqueueErr) { | 199 if (CL_SUCCESS != enqueueErr) { |
177 SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); | 200 SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); |
178 fIsGood = false; | 201 fIsGood = false; |
179 return -1; | 202 return -1; |
180 } | 203 } |
181 | 204 |
182 // This makes things totally synchronous. Actual queue is not ready yet | 205 // This makes things totally synchronous. Actual queue is not ready yet |
183 clWaitForEvents(1, &event); | 206 clWaitForEvents(1, &event); |
184 diff->finished = true; | 207 diff->finished = true; |
185 | 208 |
186 // Immediate read back the results | 209 // Immediate read back the results |
187 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, sizeof(i
nt), &numDiffPixels, 0, NULL, NULL); | 210 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, |
188 diff->result *= (double)numDiffPixels; | 211 sizeof(int), &diff->numDiffPixels, 0, NULL, NULL); |
| 212 diff->result *= (double)diff->numDiffPixels; |
189 diff->result = (1.0 - diff->result); | 213 diff->result = (1.0 - diff->result); |
190 | 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 |
191 // Release all the buffers created | 219 // Release all the buffers created |
| 220 clReleaseMemObject(diff->poiBuffer); |
192 clReleaseMemObject(diff->resultsBuffer); | 221 clReleaseMemObject(diff->resultsBuffer); |
193 clReleaseMemObject(diff->baseline); | 222 clReleaseMemObject(diff->baseline); |
194 clReleaseMemObject(diff->test); | 223 clReleaseMemObject(diff->test); |
195 | 224 |
196 SkDebugf("Time: %f\n", (get_seconds() - startTime)); | 225 SkDebugf("Time: %f\n", (get_seconds() - startTime)); |
197 | 226 |
198 return diffID; | 227 return diffID; |
199 } | 228 } |
200 | 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 |
201 bool SkDifferentPixelsImageDiffer::isFinished(int id) { | 238 bool SkDifferentPixelsImageDiffer::isFinished(int id) { |
202 return fQueuedDiffs[id].finished; | 239 return fQueuedDiffs[id].finished; |
203 } | 240 } |
204 | 241 |
205 double SkDifferentPixelsImageDiffer::getResult(int id) { | 242 double SkDifferentPixelsImageDiffer::getResult(int id) { |
206 return fQueuedDiffs[id].result; | 243 return fQueuedDiffs[id].result; |
207 } | 244 } |
208 | 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 } |
209 | 253 |
210 bool SkDifferentPixelsImageDiffer::onInit() { | 254 bool SkDifferentPixelsImageDiffer::onInit() { |
211 if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel)
) { | 255 if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel)
) { |
212 return false; | 256 return false; |
213 } | 257 } |
214 | 258 |
215 return true; | 259 return true; |
216 } | 260 } |
OLD | NEW |