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 |
157 int numDiffPixels = 0; | 173 int zero = 0; |
bsalomon
2013/06/28 14:47:17
?
Zach Reizner
2013/06/28 14:52:58
I have to initialize OpenCL buffers with pointers.
bsalomon
2013/06/28 14:56:37
ok, how about
static const int kZero;
| |
158 diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_CO PY_HOST_PTR, | 174 diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_CO PY_HOST_PTR, |
159 sizeof(int), &numDiffPixels, NULL); | 175 sizeof(int), &zero, NULL); |
176 | |
177 diff->poiBuffer = clCreateBuffer(fContext, CL_MEM_WRITE_ONLY, | |
178 sizeof(int) * 2 * baseline->width() * basel ine->height(), | |
179 NULL, NULL); | |
160 | 180 |
161 // Set all kernel arguments | 181 // Set all kernel arguments |
162 cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baselin e); | 182 cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baselin e); |
163 setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); | 183 setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); |
164 setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->results Buffer); | 184 setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->results Buffer); |
185 setArgErr |= clSetKernelArg(fKernel, 3, sizeof(cl_mem), &diff->poiBuff er); | |
165 if (CL_SUCCESS != setArgErr) { | 186 if (CL_SUCCESS != setArgErr) { |
166 SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); | 187 SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); |
167 fIsGood = false; | 188 fIsGood = false; |
168 return -1; | 189 return -1; |
169 } | 190 } |
170 | 191 |
171 // Queue this diff on the CL device | 192 // Queue this diff on the CL device |
172 cl_event event; | 193 cl_event event; |
173 const size_t workSize[] = { baseline->width(), baseline->height() }; | 194 const size_t workSize[] = { baseline->width(), baseline->height() }; |
174 cl_int enqueueErr; | 195 cl_int enqueueErr; |
175 enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSiz e, NULL, 0, NULL, &event); | 196 enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSiz e, |
197 NULL, 0, NULL, &event); | |
176 if (CL_SUCCESS != enqueueErr) { | 198 if (CL_SUCCESS != enqueueErr) { |
177 SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); | 199 SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); |
178 fIsGood = false; | 200 fIsGood = false; |
179 return -1; | 201 return -1; |
180 } | 202 } |
181 | 203 |
182 // This makes things totally synchronous. Actual queue is not ready yet | 204 // This makes things totally synchronous. Actual queue is not ready yet |
183 clWaitForEvents(1, &event); | 205 clWaitForEvents(1, &event); |
184 diff->finished = true; | 206 diff->finished = true; |
185 | 207 |
186 // Immediate read back the results | 208 // Immediate read back the results |
187 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, sizeof(i nt), &numDiffPixels, 0, NULL, NULL); | 209 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, |
188 diff->result *= (double)numDiffPixels; | 210 sizeof(int), &diff->numDiffPixels, 0, NULL, NULL); |
211 diff->result *= (double)diff->numDiffPixels; | |
189 diff->result = (1.0 - diff->result); | 212 diff->result = (1.0 - diff->result); |
190 | 213 |
214 diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels); | |
215 clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0, | |
216 sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NU LL, NULL); | |
217 | |
191 // Release all the buffers created | 218 // Release all the buffers created |
219 clReleaseMemObject(diff->poiBuffer); | |
192 clReleaseMemObject(diff->resultsBuffer); | 220 clReleaseMemObject(diff->resultsBuffer); |
193 clReleaseMemObject(diff->baseline); | 221 clReleaseMemObject(diff->baseline); |
194 clReleaseMemObject(diff->test); | 222 clReleaseMemObject(diff->test); |
195 | 223 |
196 SkDebugf("Time: %f\n", (get_seconds() - startTime)); | 224 SkDebugf("Time: %f\n", (get_seconds() - startTime)); |
197 | 225 |
198 return diffID; | 226 return diffID; |
199 } | 227 } |
200 | 228 |
229 void SkDifferentPixelsImageDiffer::deleteDiff(int id) { | |
230 QueuedDiff* diff = &fQueuedDiffs[id]; | |
231 if (NULL != diff->poi) { | |
232 SkDELETE_ARRAY(diff->poi); | |
233 diff->poi = NULL; | |
234 } | |
235 } | |
236 | |
201 bool SkDifferentPixelsImageDiffer::isFinished(int id) { | 237 bool SkDifferentPixelsImageDiffer::isFinished(int id) { |
202 return fQueuedDiffs[id].finished; | 238 return fQueuedDiffs[id].finished; |
203 } | 239 } |
204 | 240 |
205 double SkDifferentPixelsImageDiffer::getResult(int id) { | 241 double SkDifferentPixelsImageDiffer::getResult(int id) { |
206 return fQueuedDiffs[id].result; | 242 return fQueuedDiffs[id].result; |
207 } | 243 } |
208 | 244 |
245 int SkDifferentPixelsImageDiffer::getPointsOfInterestCount(int id) { | |
246 return fQueuedDiffs[id].numDiffPixels; | |
247 } | |
248 | |
249 SkIPoint* SkDifferentPixelsImageDiffer::getPointsOfInterest(int id) { | |
250 return fQueuedDiffs[id].poi; | |
251 } | |
209 | 252 |
210 bool SkDifferentPixelsImageDiffer::onInit() { | 253 bool SkDifferentPixelsImageDiffer::onInit() { |
211 if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel) ) { | 254 if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel) ) { |
212 return false; | 255 return false; |
213 } | 256 } |
214 | 257 |
215 return true; | 258 return true; |
216 } | 259 } |
OLD | NEW |