| 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 |